[Mlir-commits] [mlir] [MLIR][LLVM] Remove typed pointers from the LLVM dialect (PR #71285)

Christian Ulmann llvmlistbot at llvm.org
Sat Nov 4 11:46:16 PDT 2023


https://github.com/Dinistro updated https://github.com/llvm/llvm-project/pull/71285

>From 00200489a3f2166eef7f241a6ab7b1d42813d005 Mon Sep 17 00:00:00 2001
From: Christian Ulmann <christianulmann at gmail.com>
Date: Sat, 4 Nov 2023 15:20:30 +0100
Subject: [PATCH 1/2] [MLIR][LLVM] Remove typed pointers from the LLVM dialect

This commit removes the support for typed pointers from the LLVM
dialect. Typed pointers have been deprecated for a while and thus this
removal was announced in a PSA:
https://discourse.llvm.org/t/psa-removal-of-typed-pointers-from-the-llvm-dialect/74502

This change includes:
- Changing the ` LLVMPointerType`
- Removing remaining usages of the builders and the now removed element
  type
- Fixing assembly formats that require fully qualified pointer types
- Updating ODS pointer constraints
---
 .../include/mlir/Dialect/LLVMIR/LLVMDialect.h |   3 +-
 .../mlir/Dialect/LLVMIR/LLVMIntrinsicOps.td   |  28 +--
 .../include/mlir/Dialect/LLVMIR/LLVMOpBase.td |  28 +--
 mlir/include/mlir/Dialect/LLVMIR/LLVMOps.td   |  56 ++---
 mlir/include/mlir/Dialect/LLVMIR/LLVMTypes.td |  21 +-
 mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td   |  54 ++--
 .../Conversion/LLVMCommon/TypeConverter.cpp   |   8 +-
 mlir/lib/Dialect/LLVMIR/IR/LLVMDialect.cpp    | 233 +-----------------
 mlir/lib/Dialect/LLVMIR/IR/LLVMMemorySlot.cpp |  18 +-
 mlir/lib/Dialect/LLVMIR/IR/LLVMTypes.cpp      |  84 +------
 .../LLVMIR/Transforms/TypeConsistency.cpp     |   8 -
 mlir/lib/Dialect/OpenACC/IR/OpenACC.cpp       |  41 ++-
 mlir/lib/Dialect/OpenMP/IR/OpenMPDialect.cpp  |  27 +-
 mlir/test/Dialect/LLVMIR/global.mlir          |   5 -
 mlir/test/Dialect/LLVMIR/invalid.mlir         |  64 +----
 .../unittests/Dialect/LLVMIR/LLVMTypeTest.cpp |  34 ---
 16 files changed, 118 insertions(+), 594 deletions(-)

diff --git a/mlir/include/mlir/Dialect/LLVMIR/LLVMDialect.h b/mlir/include/mlir/Dialect/LLVMIR/LLVMDialect.h
index bbed1ea5cf62204..06df4a601b7a3f7 100644
--- a/mlir/include/mlir/Dialect/LLVMIR/LLVMDialect.h
+++ b/mlir/include/mlir/Dialect/LLVMIR/LLVMDialect.h
@@ -209,8 +209,7 @@ class GEPIndicesAdaptor {
 /// global and use it to compute the address of the first character in the
 /// string (operations inserted at the builder insertion point).
 Value createGlobalString(Location loc, OpBuilder &builder, StringRef name,
-                         StringRef value, Linkage linkage,
-                         bool useOpaquePointers = true);
+                         StringRef value, Linkage linkage);
 
 /// LLVM requires some operations to be inside of a Module operation. This
 /// function confirms that the Operation has the desired properties.
diff --git a/mlir/include/mlir/Dialect/LLVMIR/LLVMIntrinsicOps.td b/mlir/include/mlir/Dialect/LLVMIR/LLVMIntrinsicOps.td
index 72c932ac07a2e1a..1123466b7a75e3f 100644
--- a/mlir/include/mlir/Dialect/LLVMIR/LLVMIntrinsicOps.td
+++ b/mlir/include/mlir/Dialect/LLVMIR/LLVMIntrinsicOps.td
@@ -469,16 +469,16 @@ def LLVM_ThreadlocalAddressOp : LLVM_OneResultIntrOp<"threadlocal.address", [],
 
 def LLVM_CoroIdOp : LLVM_IntrOp<"coro.id", [], [], [], 1> {
   let arguments = (ins I32:$align,
-                       LLVM_i8Ptr:$promise,
-                       LLVM_i8Ptr:$coroaddr,
-                       LLVM_i8Ptr:$fnaddrs);
+                       LLVM_AnyPointer:$promise,
+                       LLVM_AnyPointer:$coroaddr,
+                       LLVM_AnyPointer:$fnaddrs);
   let assemblyFormat = "$align `,` $promise `,` $coroaddr `,` $fnaddrs"
     " attr-dict `:` functional-type(operands, results)";
 }
 
 def LLVM_CoroBeginOp : LLVM_IntrOp<"coro.begin", [], [], [], 1> {
   let arguments = (ins LLVM_TokenType:$token,
-                       LLVM_i8Ptr:$mem);
+                       LLVM_AnyPointer:$mem);
   let assemblyFormat = "$token `,` $mem attr-dict `:` functional-type(operands, results)";
 }
 
@@ -491,7 +491,7 @@ def LLVM_CoroAlignOp : LLVM_IntrOp<"coro.align", [0], [], [], 1> {
 }
 
 def LLVM_CoroSaveOp : LLVM_IntrOp<"coro.save", [], [], [], 1> {
-  let arguments = (ins LLVM_i8Ptr:$handle);
+  let arguments = (ins LLVM_AnyPointer:$handle);
   let assemblyFormat = "$handle attr-dict `:` functional-type(operands, results)";
 }
 
@@ -502,7 +502,7 @@ def LLVM_CoroSuspendOp : LLVM_IntrOp<"coro.suspend", [], [], [], 1> {
 }
 
 def LLVM_CoroEndOp : LLVM_IntrOp<"coro.end", [], [], [], 1> {
-  let arguments = (ins LLVM_i8Ptr:$handle,
+  let arguments = (ins LLVM_AnyPointer:$handle,
                        I1:$unwind,
                        LLVM_TokenType:$retvals);
   let assemblyFormat = "$handle `,` $unwind `,` $retvals attr-dict `:` functional-type(operands, results)";
@@ -510,12 +510,12 @@ def LLVM_CoroEndOp : LLVM_IntrOp<"coro.end", [], [], [], 1> {
 
 def LLVM_CoroFreeOp : LLVM_IntrOp<"coro.free", [], [], [], 1> {
   let arguments = (ins LLVM_TokenType:$id,
-                       LLVM_i8Ptr:$handle);
+                       LLVM_AnyPointer:$handle);
   let assemblyFormat = "$id `,` $handle attr-dict `:` functional-type(operands, results)";
 }
 
 def LLVM_CoroResumeOp : LLVM_IntrOp<"coro.resume", [], [], [], 0> {
-  let arguments = (ins LLVM_i8Ptr:$handle);
+  let arguments = (ins LLVM_AnyPointer:$handle);
   let assemblyFormat = "$handle attr-dict `:` qualified(type($handle))";
 }
 
@@ -591,19 +591,19 @@ def LLVM_DbgLabelOp : LLVM_IntrOp<"dbg.label", [], [], [], 0> {
 //
 
 def LLVM_VaStartOp : LLVM_ZeroResultIntrOp<"vastart">,
-                     Arguments<(ins LLVM_i8Ptr:$arg_list)> {
+                     Arguments<(ins LLVM_AnyPointer:$arg_list)> {
   let assemblyFormat = "$arg_list attr-dict `:` qualified(type($arg_list))";
   let summary = "Initializes `arg_list` for subsequent variadic argument extractions.";
 }
 
 def LLVM_VaCopyOp : LLVM_ZeroResultIntrOp<"vacopy">,
-                    Arguments<(ins LLVM_i8Ptr:$dest_list, LLVM_i8Ptr:$src_list)> {
+                    Arguments<(ins LLVM_AnyPointer:$dest_list, LLVM_AnyPointer:$src_list)> {
   let assemblyFormat = "$src_list `to` $dest_list attr-dict `:` type(operands)";
   let summary = "Copies the current argument position from `src_list` to `dest_list`.";
 }
 
 def LLVM_VaEndOp : LLVM_ZeroResultIntrOp<"vaend">,
-                   Arguments<(ins LLVM_i8Ptr:$arg_list)> {
+                   Arguments<(ins LLVM_AnyPointer:$arg_list)> {
   let assemblyFormat = "$arg_list attr-dict `:` qualified(type($arg_list))";
   let summary = "Destroys `arg_list`, which has been initialized by `intr.vastart` or `intr.vacopy`.";
 }
@@ -613,7 +613,7 @@ def LLVM_VaEndOp : LLVM_ZeroResultIntrOp<"vaend">,
 //
 
 def LLVM_EhTypeidForOp : LLVM_OneResultIntrOp<"eh.typeid.for"> {
-    let arguments = (ins LLVM_i8Ptr:$type_info);
+    let arguments = (ins LLVM_AnyPointer:$type_info);
     let assemblyFormat = "$type_info attr-dict `:` functional-type(operands, results)";
 }
 
@@ -927,12 +927,12 @@ def LLVM_PtrAnnotation
     : LLVM_OneResultIntrOp<"ptr.annotation", [0], [2],
         [AllTypesMatch<["res", "ptr"]>,
          AllTypesMatch<["annotation", "fileName", "attr"]>]> {
-  let arguments = (ins LLVM_PointerTo<AnySignlessInteger>:$ptr,
+  let arguments = (ins LLVM_AnyPointer:$ptr,
                        LLVM_AnyPointer:$annotation,
                        LLVM_AnyPointer:$fileName,
                        I32:$line,
                        LLVM_AnyPointer:$attr);
-  let results = (outs LLVM_PointerTo<AnySignlessInteger>:$res);
+  let results = (outs LLVM_AnyPointer:$res);
 }
 
 def LLVM_Annotation
diff --git a/mlir/include/mlir/Dialect/LLVMIR/LLVMOpBase.td b/mlir/include/mlir/Dialect/LLVMIR/LLVMOpBase.td
index 503988abfc090a2..0166fbb647b41b5 100644
--- a/mlir/include/mlir/Dialect/LLVMIR/LLVMOpBase.td
+++ b/mlir/include/mlir/Dialect/LLVMIR/LLVMOpBase.td
@@ -55,43 +55,17 @@ def LLVM_AnyFloat : Type<
 def LLVM_AnyPointer : Type<CPred<"::llvm::isa<::mlir::LLVM::LLVMPointerType>($_self)">,
                           "LLVM pointer type", "::mlir::LLVM::LLVMPointerType">;
 
-def LLVM_OpaquePointer : Type<
-  And<[LLVM_AnyPointer.predicate,
-       CPred<"::llvm::cast<::mlir::LLVM::LLVMPointerType>($_self).isOpaque()">]>,
-  "LLVM opaque pointer", "::mlir::LLVM::LLVMPointerType">;
-
-// Type constraint accepting LLVM pointer type with an additional constraint
-// on the element type.
-class LLVM_PointerTo<Type pointee> : Type<
-  And<[LLVM_AnyPointer.predicate,
-       Or<[LLVM_OpaquePointer.predicate,
-           SubstLeaves<
-             "$_self",
-             "::llvm::cast<::mlir::LLVM::LLVMPointerType>($_self).getElementType()",
-             pointee.predicate>]>]>,
-  "LLVM pointer to " # pointee.summary, "::mlir::LLVM::LLVMPointerType">;
-
 // Opaque pointer in a given address space.
 class LLVM_OpaquePointerInAddressSpace<int addressSpace> : Type<
-  And<[LLVM_OpaquePointer.predicate,
         CPred<
           "::llvm::cast<::mlir::LLVM::LLVMPointerType>($_self).getAddressSpace() == "
-          # addressSpace>]>,
+          # addressSpace>,
   "Opaque LLVM pointer in address space " # addressSpace,
   "::mlir::LLVM::LLVMPointerType"> {
   let builderCall = "$_builder.getType<::mlir::LLVM::LLVMPointerType>("
     # addressSpace # ")";
 }
 
-// Type constraints accepting LLVM pointer type to integer of a specific width.
-class LLVM_IntPtrBase<int width, int addressSpace = 0> : Type<
-  And<[LLVM_PointerTo<I<width>>.predicate,
-       CPred<"::llvm::cast<::mlir::LLVM::LLVMPointerType>($_self).getAddressSpace()"
-             " == " # addressSpace>]>,
-  "LLVM pointer to " # I<width>.summary>;
-
-def LLVM_i8Ptr : LLVM_IntPtrBase<8>;
-
 // Type constraint accepting any LLVM structure type.
 def LLVM_AnyStruct : Type<CPred<"::llvm::isa<::mlir::LLVM::LLVMStructType>($_self)">,
                          "LLVM structure type">;
diff --git a/mlir/include/mlir/Dialect/LLVMIR/LLVMOps.td b/mlir/include/mlir/Dialect/LLVMIR/LLVMOps.td
index 638c31b39682ea6..e12bd0ee8d9d471 100644
--- a/mlir/include/mlir/Dialect/LLVMIR/LLVMOps.td
+++ b/mlir/include/mlir/Dialect/LLVMIR/LLVMOps.td
@@ -184,16 +184,14 @@ def LLVM_AllocaOp : LLVM_Op<"alloca",
   LLVM_MemOpPatterns {
   let arguments = (ins AnyInteger:$arraySize,
                    OptionalAttr<I64Attr>:$alignment,
-                   OptionalAttr<TypeAttr>:$elem_type,
+                   TypeAttr:$elem_type,
                    UnitAttr:$inalloca);
   let results = (outs Res<LLVM_AnyPointer, "",
                           [MemAlloc<AutomaticAllocationScopeResource>]>:$res);
   string llvmInstName = "Alloca";
   string llvmBuilder = [{
     auto addrSpace = $_resultType->getPointerAddressSpace();
-    llvm::Type *elementType = moduleTranslation.convertType(
-        $elem_type ? *$elem_type
-                   : ::llvm::cast<LLVMPointerType>(op.getType()).getElementType());
+    llvm::Type *elementType = moduleTranslation.convertType($elem_type);
     auto *inst = builder.CreateAlloca(elementType, addrSpace, $arraySize);
     }] # setAlignmentCode # [{
     inst->setUsedWithInAlloca($inalloca);
@@ -207,31 +205,16 @@ def LLVM_AllocaOp : LLVM_Op<"alloca",
     $res = $_builder.create<LLVM::AllocaOp>(
       $_location, $_resultType, $arraySize,
       alignment == 0 ? IntegerAttr() : $_builder.getI64IntegerAttr(alignment),
-      TypeAttr::get(allocatedType), allocaInst->isUsedWithInAlloca());
+      allocatedType, allocaInst->isUsedWithInAlloca());
   }];
   let builders = [
-    DeprecatedOpBuilder<"the usage of typed pointers is deprecated",
-               (ins "Type":$resultType, "Value":$arraySize,
-               "unsigned":$alignment),
-    [{
-      assert(!::llvm::cast<LLVMPointerType>(resultType).isOpaque() &&
-             "pass the allocated type explicitly if opaque pointers are used");
-      if (alignment == 0)
-        return build($_builder, $_state, resultType, arraySize, IntegerAttr(),
-                     TypeAttr(), false);
-      build($_builder, $_state, resultType, arraySize,
-        $_builder.getI64IntegerAttr(alignment), TypeAttr(), false);
-    }]>,
     OpBuilder<(ins "Type":$resultType, "Type":$elementType, "Value":$arraySize,
                CArg<"unsigned", "0">:$alignment),
     [{
-      TypeAttr elemTypeAttr =
-          ::llvm::cast<LLVMPointerType>(resultType).isOpaque() ?
-          TypeAttr::get(elementType) : TypeAttr();
       build($_builder, $_state, resultType, arraySize,
             alignment == 0 ? IntegerAttr()
                            : $_builder.getI64IntegerAttr(alignment),
-            elemTypeAttr, false);
+            elementType, false);
 
     }]>
     ];
@@ -247,7 +230,7 @@ def LLVM_GEPOp : LLVM_Op<"getelementptr", [Pure,
   let arguments = (ins LLVM_ScalarOrVectorOf<LLVM_AnyPointer>:$base,
                    Variadic<LLVM_ScalarOrVectorOf<AnyInteger>>:$dynamicIndices,
                    DenseI32ArrayAttr:$rawConstantIndices,
-                   OptionalAttr<TypeAttr>:$elem_type,
+                   TypeAttr:$elem_type,
                    UnitAttr:$inbounds);
   let results = (outs LLVM_ScalarOrVectorOf<LLVM_AnyPointer>:$res);
   let skipDefaultBuilders = 1;
@@ -282,14 +265,6 @@ def LLVM_GEPOp : LLVM_Op<"getelementptr", [Pure,
     OpBuilder<(ins "Type":$resultType, "Type":$basePtrType, "Value":$basePtr,
                "ValueRange":$indices, CArg<"bool", "false">:$inbounds,
                CArg<"ArrayRef<NamedAttribute>", "{}">:$attributes)>,
-    DeprecatedOpBuilder<"the usage of typed pointers is deprecated",
-               (ins "Type":$resultType, "Value":$basePtr,
-               "ValueRange":$indices, CArg<"bool", "false">:$inbounds,
-               CArg<"ArrayRef<NamedAttribute>", "{}">:$attributes)>,
-    DeprecatedOpBuilder<"the usage of typed pointers is deprecated",
-               (ins "Type":$resultType, "Value":$basePtr,
-               "ArrayRef<GEPArg>":$indices, CArg<"bool", "false">:$inbounds,
-               CArg<"ArrayRef<NamedAttribute>", "{}">:$attributes)>,
     OpBuilder<(ins "Type":$resultType, "Type":$basePtrType, "Value":$basePtr,
                "ArrayRef<GEPArg>":$indices, CArg<"bool", "false">:$inbounds,
                CArg<"ArrayRef<NamedAttribute>", "{}">:$attributes)>,
@@ -313,7 +288,7 @@ def LLVM_GEPOp : LLVM_Op<"getelementptr", [Pure,
   let assemblyFormat = [{
     (`inbounds` $inbounds^)?
     $base `[` custom<GEPIndices>($dynamicIndices, $rawConstantIndices) `]` attr-dict
-    `:` functional-type(operands, results) (`,` $elem_type^)?
+    `:` functional-type(operands, results) `,` $elem_type
   }];
 
   let extraClassDeclaration = [{
@@ -332,7 +307,7 @@ def LLVM_LoadOp : LLVM_MemAccessOpBase<"load",
     [DeclareOpInterfaceMethods<MemoryEffectsOpInterface>,
      DeclareOpInterfaceMethods<PromotableMemOpInterface>,
      DeclareOpInterfaceMethods<SafeMemorySlotAccessOpInterface>]> {
-  dag args = (ins LLVM_PointerTo<LLVM_LoadableType>:$addr,
+  dag args = (ins LLVM_AnyPointer:$addr,
               OptionalAttr<I64Attr>:$alignment,
               UnitAttr:$volatile_,
               UnitAttr:$nontemporal,
@@ -370,7 +345,8 @@ def LLVM_LoadOp : LLVM_MemAccessOpBase<"load",
   let assemblyFormat = [{
     (`volatile` $volatile_^)? $addr
     (`atomic` (`syncscope` `(` $syncscope^ `)`)? $ordering^)?
-    attr-dict `:` custom<LoadType>(type($addr), type($res))
+    attr-dict `:` qualified(type($addr)) `->` type($res)
+
   }];
   string llvmBuilder = [{
     auto *inst = builder.CreateLoad($_resultType, $addr, $volatile_);
@@ -391,9 +367,6 @@ def LLVM_LoadOp : LLVM_MemAccessOpBase<"load",
         getLLVMSyncScope(loadInst));
   }];
   let builders = [
-    DeprecatedOpBuilder<"the usage of typed pointers is deprecated",
-      (ins "Value":$addr, CArg<"unsigned", "0">:$alignment,
-      CArg<"bool", "false">:$isVolatile, CArg<"bool", "false">:$isNonTemporal)>,
     OpBuilder<(ins "Type":$type, "Value":$addr,
       CArg<"unsigned", "0">:$alignment, CArg<"bool", "false">:$isVolatile,
       CArg<"bool", "false">:$isNonTemporal,
@@ -408,7 +381,7 @@ def LLVM_StoreOp : LLVM_MemAccessOpBase<"store",
      DeclareOpInterfaceMethods<PromotableMemOpInterface>,
      DeclareOpInterfaceMethods<SafeMemorySlotAccessOpInterface>]> {
   dag args = (ins LLVM_LoadableType:$value,
-              LLVM_PointerTo<LLVM_LoadableType>:$addr,
+              LLVM_AnyPointer:$addr,
               OptionalAttr<I64Attr>:$alignment,
               UnitAttr:$volatile_,
               UnitAttr:$nontemporal,
@@ -445,7 +418,7 @@ def LLVM_StoreOp : LLVM_MemAccessOpBase<"store",
   let assemblyFormat = [{
     (`volatile` $volatile_^)? $value `,` $addr
     (`atomic` (`syncscope` `(` $syncscope^ `)`)? $ordering^)?
-    attr-dict `:` custom<StoreType>(type($value), type($addr))
+    attr-dict `:` type($value) `,` qualified(type($addr))
   }];
   string llvmBuilder = [{
     auto *inst = builder.CreateStore($value, $addr, $volatile_);
@@ -651,8 +624,7 @@ def LLVM_CallOp : LLVM_MemAccessOpBase<"call",
     OpBuilder<(ins "LLVMFunctionType":$calleeType, "FlatSymbolRefAttr":$callee,
                    CArg<"ValueRange", "{}">:$args)>,
     OpBuilder<(ins "LLVMFunctionType":$calleeType, "StringRef":$callee,
-                   CArg<"ValueRange", "{}">:$args)>,
-    OpBuilder<(ins "Value":$callee, "ValueRange":$args)>
+                   CArg<"ValueRange", "{}">:$args)>
   ];
   let hasCustomAssemblyFormat = 1;
   let extraClassDeclaration = [{
@@ -1636,7 +1608,7 @@ def LLVM_AtomicRMWOp : LLVM_MemAccessOpBase<"atomicrmw", [
       TypesMatchWith<"result #0 and operand #1 have the same type",
                      "val", "res", "$_self">]> {
   dag args = (ins AtomicBinOp:$bin_op,
-              LLVM_PointerTo<LLVM_AtomicRMWType>:$ptr,
+              LLVM_AnyPointer:$ptr,
               LLVM_AtomicRMWType:$val, AtomicOrdering:$ordering,
               OptionalAttr<StrAttr>:$syncscope,
               OptionalAttr<I64Attr>:$alignment,
@@ -1687,7 +1659,7 @@ def LLVM_AtomicCmpXchgOp : LLVM_MemAccessOpBase<"cmpxchg", [
       TypesMatchWith<"result #0 has an LLVM struct type consisting of "
                      "the type of operand #2 and a bool", "val", "res",
                      "getValAndBoolStructType($_self)">]> {
-  dag args = (ins LLVM_PointerTo<LLVM_AtomicCmpXchgType>:$ptr,
+  dag args = (ins LLVM_AnyPointer:$ptr,
               LLVM_AtomicCmpXchgType:$cmp, LLVM_AtomicCmpXchgType:$val,
               AtomicOrdering:$success_ordering,
               AtomicOrdering:$failure_ordering,
diff --git a/mlir/include/mlir/Dialect/LLVMIR/LLVMTypes.td b/mlir/include/mlir/Dialect/LLVMIR/LLVMTypes.td
index 5f746e21e831b59..f6ed05640a5fff1 100644
--- a/mlir/include/mlir/Dialect/LLVMIR/LLVMTypes.td
+++ b/mlir/include/mlir/Dialect/LLVMIR/LLVMTypes.td
@@ -139,30 +139,17 @@ def LLVMPointerType : LLVMType<"LLVMPointer", "ptr", [
     ```
   }];
 
-  let parameters = (ins DefaultValuedParameter<"Type", "Type()">:$elementType,
-                        DefaultValuedParameter<"unsigned", "0">:$addressSpace);
+  let parameters = (ins DefaultValuedParameter<"unsigned", "0">:$addressSpace);
   let assemblyFormat = [{
-    (`<` custom<Pointer>($elementType, $addressSpace)^ `>`)?
+    (`<` $addressSpace^ `>`)?
   }];
 
-  let genVerifyDecl = 1;
-
+  let skipDefaultBuilders = 1;
   let builders = [
-    TypeBuilderWithInferredContext<(ins "Type":$elementType,
-                                         CArg<"unsigned", "0">:$addressSpace)>,
     TypeBuilder<(ins CArg<"unsigned", "0">:$addressSpace), [{
-      return $_get($_ctxt, Type(), addressSpace);
+      return $_get($_ctxt, 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);
-  }];
 }
 
 //===----------------------------------------------------------------------===//
diff --git a/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td b/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
index 6947cf10e3600d4..406519d3eaa7272 100644
--- a/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
+++ b/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
@@ -19,10 +19,8 @@ include "mlir/Dialect/LLVMIR/LLVMOpBase.td"
 include "mlir/Interfaces/SideEffectInterfaces.td"
 include "mlir/Dialect/LLVMIR/BasicPtxBuilderInterface.td"
 
-def LLVM_i8Ptr_global : LLVM_IntPtrBase<8, 1>;
-def LLVM_i8Ptr_shared : LLVM_IntPtrBase<8, 3>;
-def LLVM_i64ptr_any : LLVM_IntPtrBase<64>;
-def LLVM_i64ptr_shared : LLVM_IntPtrBase<64, 3>;
+def LLVM_ptr_global : LLVM_OpaquePointerInAddressSpace<1>;
+def LLVM_ptr_shared : LLVM_OpaquePointerInAddressSpace<3>;
 
 //===----------------------------------------------------------------------===//
 // NVVM dialect definitions
@@ -213,7 +211,7 @@ def NVVM_ReduxOp :
 
 /// mbarrier.init instruction with generic pointer type
 def NVVM_MBarrierInitOp : NVVM_PTXBuilder_Op<"mbarrier.init">,
-  Arguments<(ins LLVM_i64ptr_any:$addr, I32:$count, PtxPredicate:$predicate)> {
+  Arguments<(ins LLVM_AnyPointer:$addr, I32:$count, PtxPredicate:$predicate)> {
   string llvmBuilder = [{
       createIntrinsicCall(builder, llvm::Intrinsic::nvvm_mbarrier_init, {$addr, $count});
   }];
@@ -228,7 +226,7 @@ def NVVM_MBarrierInitOp : NVVM_PTXBuilder_Op<"mbarrier.init">,
 
 /// mbarrier.init instruction with shared pointer type
 def NVVM_MBarrierInitSharedOp : NVVM_PTXBuilder_Op<"mbarrier.init.shared">,
-  Arguments<(ins LLVM_i64ptr_shared:$addr, I32:$count, PtxPredicate:$predicate)> {
+  Arguments<(ins LLVM_ptr_shared:$addr, I32:$count, PtxPredicate:$predicate)> {
   string llvmBuilder = [{
       createIntrinsicCall(builder, llvm::Intrinsic::nvvm_mbarrier_init_shared, {$addr, $count});
   }];
@@ -240,7 +238,7 @@ def NVVM_MBarrierInitSharedOp : NVVM_PTXBuilder_Op<"mbarrier.init.shared">,
 }
 
 def NVVM_MBarrierInvalOp : NVVM_Op<"mbarrier.inval">,
-  Arguments<(ins LLVM_i64ptr_any:$addr)> {
+  Arguments<(ins LLVM_AnyPointer:$addr)> {
   string llvmBuilder = [{
       createIntrinsicCall(builder, llvm::Intrinsic::nvvm_mbarrier_inval, {$addr});
   }];
@@ -248,7 +246,7 @@ def NVVM_MBarrierInvalOp : NVVM_Op<"mbarrier.inval">,
 }
 
 def NVVM_MBarrierInvalSharedOp : NVVM_Op<"mbarrier.inval.shared">,
-  Arguments<(ins LLVM_i64ptr_shared:$addr)> {
+  Arguments<(ins LLVM_ptr_shared:$addr)> {
   string llvmBuilder = [{
       createIntrinsicCall(builder, llvm::Intrinsic::nvvm_mbarrier_inval_shared, {$addr});
   }];
@@ -257,7 +255,7 @@ def NVVM_MBarrierInvalSharedOp : NVVM_Op<"mbarrier.inval.shared">,
 
 def NVVM_MBarrierArriveOp : NVVM_Op<"mbarrier.arrive">,
   Results<(outs LLVM_Type:$res)>,
-  Arguments<(ins LLVM_i64ptr_any:$addr)> {
+  Arguments<(ins LLVM_AnyPointer:$addr)> {
   string llvmBuilder = [{
       $res = createIntrinsicCall(builder, llvm::Intrinsic::nvvm_mbarrier_arrive, {$addr});
   }];
@@ -266,16 +264,16 @@ def NVVM_MBarrierArriveOp : NVVM_Op<"mbarrier.arrive">,
 
 def NVVM_MBarrierArriveSharedOp : NVVM_Op<"mbarrier.arrive.shared">,
   Results<(outs LLVM_Type:$res)>,
-  Arguments<(ins LLVM_i64ptr_shared:$addr)> {
+  Arguments<(ins LLVM_ptr_shared:$addr)> {
   string llvmBuilder = [{
       $res = createIntrinsicCall(builder, llvm::Intrinsic::nvvm_mbarrier_arrive_shared, {$addr});
   }];
-  let assemblyFormat = "$addr attr-dict `:` type($addr) `->` type($res)";
+  let assemblyFormat = "$addr attr-dict `:` qualified(type($addr)) `->` type($res)";
 }
 
 def NVVM_MBarrierArriveNocompleteOp : NVVM_Op<"mbarrier.arrive.nocomplete">,
   Results<(outs LLVM_Type:$res)>,
-  Arguments<(ins LLVM_i64ptr_any:$addr, I32:$count)> {
+  Arguments<(ins LLVM_AnyPointer:$addr, I32:$count)> {
   string llvmBuilder = [{
       $res = createIntrinsicCall(builder, llvm::Intrinsic::nvvm_mbarrier_arrive_noComplete, {$addr, $count});
   }];
@@ -284,7 +282,7 @@ def NVVM_MBarrierArriveNocompleteOp : NVVM_Op<"mbarrier.arrive.nocomplete">,
 
 def NVVM_MBarrierArriveNocompleteSharedOp : NVVM_Op<"mbarrier.arrive.nocomplete.shared">,
   Results<(outs LLVM_Type:$res)>,
-  Arguments<(ins LLVM_i64ptr_shared:$addr, I32:$count)> {
+  Arguments<(ins LLVM_ptr_shared:$addr, I32:$count)> {
   string llvmBuilder = [{
       $res = createIntrinsicCall(builder, llvm::Intrinsic::nvvm_mbarrier_arrive_noComplete_shared, {$addr, $count});
   }];
@@ -292,7 +290,7 @@ def NVVM_MBarrierArriveNocompleteSharedOp : NVVM_Op<"mbarrier.arrive.nocomplete.
 }
 
 def NVVM_MBarrierArriveExpectTxOp : NVVM_PTXBuilder_Op<"mbarrier.arrive.expect_tx">,  
-  Arguments<(ins LLVM_i64ptr_any:$addr, I32:$txcount, PtxPredicate:$predicate)> {
+  Arguments<(ins LLVM_AnyPointer:$addr, I32:$txcount, PtxPredicate:$predicate)> {
   let assemblyFormat = "$addr `,` $txcount (`,` `predicate` `=` $predicate^)? attr-dict `:` type(operands)";
   let extraClassDefinition = [{
     std::string $cppClass::getPtx() { return std::string("mbarrier.arrive.expect_tx.b64 _, [%0], %1;"); }
@@ -300,7 +298,7 @@ def NVVM_MBarrierArriveExpectTxOp : NVVM_PTXBuilder_Op<"mbarrier.arrive.expect_t
 }
 
 def NVVM_MBarrierArriveExpectTxSharedOp : NVVM_PTXBuilder_Op<"mbarrier.arrive.expect_tx.shared">,  
-  Arguments<(ins LLVM_i64ptr_shared:$addr, I32:$txcount, PtxPredicate:$predicate)> {    
+  Arguments<(ins LLVM_ptr_shared:$addr, I32:$txcount, PtxPredicate:$predicate)> {    
   let assemblyFormat = "$addr `,` $txcount (`,` `predicate` `=` $predicate^)? attr-dict `:` type(operands)";
   let extraClassDefinition = [{
     std::string $cppClass::getPtx() { return std::string("mbarrier.arrive.expect_tx.shared.b64 _, [%0], %1;"); }
@@ -308,7 +306,7 @@ def NVVM_MBarrierArriveExpectTxSharedOp : NVVM_PTXBuilder_Op<"mbarrier.arrive.ex
 }
 
 def NVVM_MBarrierTryWaitParityOp : NVVM_PTXBuilder_Op<"mbarrier.try_wait.parity">,  
-  Arguments<(ins LLVM_i64ptr_any:$addr, I32:$phase, I32:$ticks)> {  
+  Arguments<(ins LLVM_AnyPointer:$addr, I32:$phase, I32:$ticks)> {  
   let assemblyFormat = "$addr `,` $phase `,` $ticks attr-dict `:` type(operands)";
   let extraClassDefinition = [{
     std::string $cppClass::getPtx() {
@@ -327,7 +325,7 @@ def NVVM_MBarrierTryWaitParityOp : NVVM_PTXBuilder_Op<"mbarrier.try_wait.parity"
 }
 
 def NVVM_MBarrierTryWaitParitySharedOp : NVVM_PTXBuilder_Op<"mbarrier.try_wait.parity.shared">,  
-  Arguments<(ins LLVM_i64ptr_shared:$addr, I32:$phase, I32:$ticks)> {  
+  Arguments<(ins LLVM_ptr_shared:$addr, I32:$phase, I32:$ticks)> {  
   let assemblyFormat = "$addr `,` $phase `,` $ticks attr-dict `:` type(operands)";
   let extraClassDefinition = [{
     std::string $cppClass::getPtx() {
@@ -347,7 +345,7 @@ def NVVM_MBarrierTryWaitParitySharedOp : NVVM_PTXBuilder_Op<"mbarrier.try_wait.p
 
 def NVVM_MBarrierTestWaitOp : NVVM_Op<"mbarrier.test.wait">,
   Results<(outs LLVM_Type:$res)>,
-  Arguments<(ins LLVM_i64ptr_any:$addr, LLVM_Type:$state)> {
+  Arguments<(ins LLVM_AnyPointer:$addr, LLVM_Type:$state)> {
   string llvmBuilder = [{
       $res = createIntrinsicCall(builder, llvm::Intrinsic::nvvm_mbarrier_test_wait, {$addr, $state});
   }];
@@ -356,7 +354,7 @@ def NVVM_MBarrierTestWaitOp : NVVM_Op<"mbarrier.test.wait">,
 
 def NVVM_MBarrierTestWaitSharedOp : NVVM_Op<"mbarrier.test.wait.shared">,
   Results<(outs LLVM_Type:$res)>,
-  Arguments<(ins LLVM_i64ptr_shared:$addr, LLVM_Type:$state)> {
+  Arguments<(ins LLVM_ptr_shared:$addr, LLVM_Type:$state)> {
   string llvmBuilder = [{
       $res = createIntrinsicCall(builder, llvm::Intrinsic::nvvm_mbarrier_test_wait_shared, {$addr, $state});
   }];
@@ -501,8 +499,8 @@ def LoadCacheModifierKind : I32EnumAttr<"LoadCacheModifierKind",
 def LoadCacheModifierAttr : EnumAttr<NVVM_Dialect, LoadCacheModifierKind, "load_cache_modifier">;
 
 def NVVM_CpAsyncOp : NVVM_PTXBuilder_Op<"cp.async.shared.global">,
-  Arguments<(ins LLVM_i8Ptr_shared:$dst,
-                 LLVM_i8Ptr_global:$src,
+  Arguments<(ins LLVM_ptr_shared:$dst,
+                 LLVM_ptr_global:$src,
                  I32Attr:$size,
                  LoadCacheModifierAttr:$modifier,
                  Optional<LLVM_Type>:$cpSize)> {
@@ -1187,7 +1185,7 @@ def NVVM_WMMAMmaOp : NVVM_Op<"wmma.mma">,
 }
 
 def NVVM_StMatrixOp: NVVM_PTXBuilder_Op<"stmatrix">, 
-  Arguments<(ins LLVM_i8Ptr_shared:$ptr, 
+  Arguments<(ins LLVM_ptr_shared:$ptr, 
                  Variadic<I32>:$sources, 
                  MMALayoutAttr:$layout)> {
   let summary = "cooperative matrix store";
@@ -1404,9 +1402,9 @@ def NVVM_CpAsyncBulkTensorGlobalToSharedClusterOp :
   NVVM_Op<"cp.async.bulk.tensor.shared.cluster.global", 
   [DeclareOpInterfaceMethods<BasicPtxBuilderOpInterface>, 
   AttrSizedOperandSegments]>,
-  Arguments<(ins  LLVM_i64ptr_shared:$dstMem,
-                  LLVM_i64ptr_any:$tmaDescriptor,
-                  LLVM_i64ptr_shared:$mbar,
+  Arguments<(ins  LLVM_ptr_shared:$dstMem,
+                  LLVM_AnyPointer:$tmaDescriptor,
+                  LLVM_ptr_shared:$mbar,
                   Variadic<I32>:$coordinates,
                   PtxPredicate:$predicate)> {
   let assemblyFormat = [{ 
@@ -1439,8 +1437,8 @@ def NVVM_CpAsyncBulkTensorSharedCTAToGlobalOp :
   NVVM_Op<"cp.async.bulk.tensor.global.shared.cta", 
   [DeclareOpInterfaceMethods<BasicPtxBuilderOpInterface>, 
   AttrSizedOperandSegments]>,
-  Arguments<(ins  LLVM_i64ptr_any:$tmaDescriptor,
-                  LLVM_i64ptr_shared:$srcMem,
+  Arguments<(ins  LLVM_AnyPointer:$tmaDescriptor,
+                  LLVM_ptr_shared:$srcMem,
                   Variadic<I32>:$coordinates,
                   PtxPredicate:$predicate)> {
   let assemblyFormat = [{ 
@@ -1469,7 +1467,7 @@ def NVVM_CpAsyncBulkTensorSharedCTAToGlobalOp :
 
 def NVVM_PrefetchTensorMapOp : NVVM_Op<"prefetch.tensormap",
                     [DeclareOpInterfaceMethods<BasicPtxBuilderOpInterface>]>,
-  Arguments<(ins LLVM_i64ptr_any:$tmaDescriptor, PtxPredicate:$predicate)> {
+  Arguments<(ins LLVM_AnyPointer:$tmaDescriptor, PtxPredicate:$predicate)> {
   let assemblyFormat = "$tmaDescriptor (`,` `predicate` `=` $predicate^)? attr-dict `:` type(operands)";
   let extraClassDefinition = [{
     std::string $cppClass::getPtx() { 
diff --git a/mlir/lib/Conversion/LLVMCommon/TypeConverter.cpp b/mlir/lib/Conversion/LLVMCommon/TypeConverter.cpp
index 35b95d7a5ebe925..fdd1d993fdb8f62 100644
--- a/mlir/lib/Conversion/LLVMCommon/TypeConverter.cpp
+++ b/mlir/lib/Conversion/LLVMCommon/TypeConverter.cpp
@@ -78,13 +78,7 @@ LLVMTypeConverter::LLVMTypeConverter(MLIRContext *ctx,
 
   // LLVM container types may (recursively) contain other types that must be
   // converted even when the outer type is compatible.
-  addConversion([&](LLVM::LLVMPointerType type) -> std::optional<Type> {
-    if (type.isOpaque())
-      return type;
-    if (auto pointee = convertType(type.getElementType()))
-      return LLVM::LLVMPointerType::get(pointee, type.getAddressSpace());
-    return std::nullopt;
-  });
+  addConversion([&](LLVM::LLVMPointerType type) { return type; });
 
   addConversion([&](LLVM::LLVMStructType type, SmallVectorImpl<Type> &results)
                     -> std::optional<LogicalResult> {
diff --git a/mlir/lib/Dialect/LLVMIR/IR/LLVMDialect.cpp b/mlir/lib/Dialect/LLVMIR/IR/LLVMDialect.cpp
index 7f5681e7bdc0592..3ce7e9abccaa387 100644
--- a/mlir/lib/Dialect/LLVMIR/IR/LLVMDialect.cpp
+++ b/mlir/lib/Dialect/LLVMIR/IR/LLVMDialect.cpp
@@ -216,17 +216,13 @@ OpFoldResult ICmpOp::fold(FoldAdaptor adaptor) {
 //===----------------------------------------------------------------------===//
 
 void AllocaOp::print(OpAsmPrinter &p) {
-  Type elemTy = llvm::cast<LLVM::LLVMPointerType>(getType()).getElementType();
-  if (!elemTy)
-    elemTy = *getElemType();
-
   auto funcTy =
       FunctionType::get(getContext(), {getArraySize().getType()}, {getType()});
 
   if (getInalloca())
     p << " inalloca";
 
-  p << ' ' << getArraySize() << " x " << elemTy;
+  p << ' ' << getArraySize() << " x " << getElemType();
   if (getAlignment() && *getAlignment() != 0)
     p.printOptionalAttrDict((*this)->getAttrs(),
                             {kElemTypeAttrName, getInallocaAttrName()});
@@ -277,40 +273,16 @@ ParseResult AllocaOp::parse(OpAsmParser &parser, OperationState &result) {
     return failure();
 
   Type resultType = funcType.getResult(0);
-  if (auto ptrResultType = llvm::dyn_cast<LLVMPointerType>(resultType)) {
-    if (ptrResultType.isOpaque())
-      result.addAttribute(kElemTypeAttrName, TypeAttr::get(elemType));
-  }
+  if (auto ptrResultType = llvm::dyn_cast<LLVMPointerType>(resultType))
+    result.addAttribute(kElemTypeAttrName, TypeAttr::get(elemType));
 
   result.addTypes({funcType.getResult(0)});
   return success();
 }
 
-/// Checks that the elemental type is present in either the pointer type or
-/// the attribute, but not both.
-static LogicalResult verifyOpaquePtr(Operation *op, LLVMPointerType ptrType,
-                                     std::optional<Type> ptrElementType) {
-  if (ptrType.isOpaque() && !ptrElementType.has_value()) {
-    return op->emitOpError() << "expected '" << kElemTypeAttrName
-                             << "' attribute if opaque pointer type is used";
-  }
-  if (!ptrType.isOpaque() && ptrElementType.has_value()) {
-    return op->emitOpError()
-           << "unexpected '" << kElemTypeAttrName
-           << "' attribute when non-opaque pointer type is used";
-  }
-  return success();
-}
-
 LogicalResult AllocaOp::verify() {
-  LLVMPointerType ptrType = llvm::cast<LLVMPointerType>(getType());
-  if (failed(verifyOpaquePtr(getOperation(), ptrType, getElemType())))
-    return failure();
-
-  Type elemTy =
-      (ptrType.isOpaque()) ? *getElemType() : ptrType.getElementType();
   // Only certain target extension types can be used in 'alloca'.
-  if (auto targetExtType = dyn_cast<LLVMTargetExtType>(elemTy);
+  if (auto targetExtType = dyn_cast<LLVMTargetExtType>(getElemType());
       targetExtType && !targetExtType.supportsMemOps())
     return emitOpError()
            << "this target extension type cannot be used in alloca";
@@ -318,11 +290,7 @@ LogicalResult AllocaOp::verify() {
   return success();
 }
 
-Type AllocaOp::getResultPtrElementType() {
-  // This will become trivial once non-opaque pointers are gone.
-  return getElemType().has_value() ? *getElemType()
-                                   : getResult().getType().getElementType();
-}
+Type AllocaOp::getResultPtrElementType() { return getElemType(); }
 
 //===----------------------------------------------------------------------===//
 // LLVM::BrOp
@@ -525,18 +493,6 @@ static Type extractVectorElementType(Type type) {
   return type;
 }
 
-void GEPOp::build(OpBuilder &builder, OperationState &result, Type resultType,
-                  Value basePtr, ArrayRef<GEPArg> indices, bool inbounds,
-                  ArrayRef<NamedAttribute> attributes) {
-  auto ptrType =
-      llvm::cast<LLVMPointerType>(extractVectorElementType(basePtr.getType()));
-  assert(!ptrType.isOpaque() &&
-         "expected non-opaque pointer, provide elementType explicitly when "
-         "opaque pointers are used");
-  build(builder, result, resultType, ptrType.getElementType(), basePtr, indices,
-        inbounds, attributes);
-}
-
 /// Destructures the 'indices' parameter into 'rawConstantIndices' and
 /// 'dynamicIndices', encoding the former in the process. In the process,
 /// dynamic indices which are used to index into a structure type are converted
@@ -602,20 +558,11 @@ void GEPOp::build(OpBuilder &builder, OperationState &result, Type resultType,
     result.addAttribute(getInboundsAttrName(result.name),
                         builder.getUnitAttr());
   }
-  if (llvm::cast<LLVMPointerType>(extractVectorElementType(basePtr.getType()))
-          .isOpaque())
-    result.addAttribute(kElemTypeAttrName, TypeAttr::get(elementType));
+  result.addAttribute(kElemTypeAttrName, TypeAttr::get(elementType));
   result.addOperands(basePtr);
   result.addOperands(dynamicIndices);
 }
 
-void GEPOp::build(OpBuilder &builder, OperationState &result, Type resultType,
-                  Value basePtr, ValueRange indices, bool inbounds,
-                  ArrayRef<NamedAttribute> attributes) {
-  build(builder, result, resultType, basePtr, SmallVector<GEPArg>(indices),
-        inbounds, attributes);
-}
-
 void GEPOp::build(OpBuilder &builder, OperationState &result, Type resultType,
                   Type elementType, Value basePtr, ValueRange indices,
                   bool inbounds, ArrayRef<NamedAttribute> attributes) {
@@ -712,12 +659,6 @@ verifyStructIndices(Type baseGEPType, GEPIndicesAdaptor<ValueRange> indices,
 }
 
 LogicalResult LLVM::GEPOp::verify() {
-  if (failed(verifyOpaquePtr(
-          getOperation(),
-          llvm::cast<LLVMPointerType>(extractVectorElementType(getType())),
-          getElemType())))
-    return failure();
-
   if (static_cast<size_t>(
           llvm::count(getRawConstantIndices(), kDynamicIndex)) !=
       getDynamicIndices().size())
@@ -728,14 +669,7 @@ LogicalResult LLVM::GEPOp::verify() {
                              [&] { return emitOpError(); });
 }
 
-Type LLVM::GEPOp::getSourceElementType() {
-  if (std::optional<Type> elemType = getElemType())
-    return *elemType;
-
-  return llvm::cast<LLVMPointerType>(
-             extractVectorElementType(getBase().getType()))
-      .getElementType();
-}
+Type LLVM::GEPOp::getSourceElementType() { return getElemType(); }
 
 Type GEPOp::getResultPtrElementType() {
   // Set the initial type currently being used for indexing. This will be
@@ -836,14 +770,6 @@ LogicalResult LoadOp::verify() {
                            {AtomicOrdering::release, AtomicOrdering::acq_rel});
 }
 
-void LoadOp::build(OpBuilder &builder, OperationState &state, Value addr,
-                   unsigned alignment, bool isVolatile, bool isNonTemporal) {
-  auto type = llvm::cast<LLVMPointerType>(addr.getType()).getElementType();
-  assert(type && "must provide explicit element type to the constructor "
-                 "when the pointer type is opaque");
-  build(builder, state, type, addr, alignment, isVolatile, isNonTemporal);
-}
-
 void LoadOp::build(OpBuilder &builder, OperationState &state, Type type,
                    Value addr, unsigned alignment, bool isVolatile,
                    bool isNonTemporal, AtomicOrdering ordering,
@@ -857,51 +783,6 @@ void LoadOp::build(OpBuilder &builder, OperationState &state, Type type,
         /*tbaa=*/nullptr);
 }
 
-// Extract the pointee type from the LLVM pointer type wrapped in MLIR. Return
-// the resulting type if any, null type if opaque pointers are used, and
-// std::nullopt if the given type is not the pointer type.
-static std::optional<Type>
-getLoadStoreElementType(OpAsmParser &parser, Type type, SMLoc trailingTypeLoc) {
-  auto llvmTy = llvm::dyn_cast<LLVM::LLVMPointerType>(type);
-  if (!llvmTy) {
-    parser.emitError(trailingTypeLoc, "expected LLVM pointer type");
-    return std::nullopt;
-  }
-  return llvmTy.getElementType();
-}
-
-/// Parses the LoadOp type either using the typed or opaque pointer format.
-// TODO: Drop once the typed pointer assembly format is not needed anymore.
-static ParseResult parseLoadType(OpAsmParser &parser, Type &type,
-                                 Type &elementType) {
-  SMLoc trailingTypeLoc;
-  if (parser.getCurrentLocation(&trailingTypeLoc) || parser.parseType(type))
-    return failure();
-
-  std::optional<Type> pointerElementType =
-      getLoadStoreElementType(parser, type, trailingTypeLoc);
-  if (!pointerElementType)
-    return failure();
-  if (*pointerElementType) {
-    elementType = *pointerElementType;
-    return success();
-  }
-
-  if (parser.parseArrow() || parser.parseType(elementType))
-    return failure();
-  return success();
-}
-
-/// Prints the LoadOp type either using the typed or opaque pointer format.
-// TODO: Drop once the typed pointer assembly format is not needed anymore.
-static void printLoadType(OpAsmPrinter &printer, Operation *op, Type type,
-                          Type elementType) {
-  printer << type;
-  auto pointerType = cast<LLVMPointerType>(type);
-  if (pointerType.isOpaque())
-    printer << " -> " << elementType;
-}
-
 //===----------------------------------------------------------------------===//
 // StoreOp
 //===----------------------------------------------------------------------===//
@@ -940,38 +821,6 @@ void StoreOp::build(OpBuilder &builder, OperationState &state, Value value,
         /*alias_scopes=*/nullptr, /*noalias_scopes=*/nullptr, /*tbaa=*/nullptr);
 }
 
-/// Parses the StoreOp type either using the typed or opaque pointer format.
-// TODO: Drop once the typed pointer assembly format is not needed anymore.
-static ParseResult parseStoreType(OpAsmParser &parser, Type &elementType,
-                                  Type &type) {
-  SMLoc trailingTypeLoc;
-  if (parser.getCurrentLocation(&trailingTypeLoc) ||
-      parser.parseType(elementType))
-    return failure();
-
-  if (succeeded(parser.parseOptionalComma()))
-    return parser.parseType(type);
-
-  // Extract the element type from the pointer type.
-  type = elementType;
-  std::optional<Type> pointerElementType =
-      getLoadStoreElementType(parser, type, trailingTypeLoc);
-  if (!pointerElementType)
-    return failure();
-  elementType = *pointerElementType;
-  return success();
-}
-
-/// Prints the StoreOp type either using the typed or opaque pointer format.
-// TODO: Drop once the typed pointer assembly format is not needed anymore.
-static void printStoreType(OpAsmPrinter &printer, Operation *op,
-                           Type elementType, Type type) {
-  auto pointerType = cast<LLVMPointerType>(type);
-  if (pointerType.isOpaque())
-    printer << elementType << ", ";
-  printer << type;
-}
-
 //===----------------------------------------------------------------------===//
 // CallOp
 //===----------------------------------------------------------------------===//
@@ -1055,22 +904,6 @@ void CallOp::build(OpBuilder &builder, OperationState &state, LLVMFuncOp func,
         /*access_groups=*/nullptr, /*alias_scopes=*/nullptr,
         /*noalias_scopes=*/nullptr, /*tbaa=*/nullptr);
 }
-
-void CallOp::build(OpBuilder &builder, OperationState &state, Value callee,
-                   ValueRange args) {
-  auto calleeType = cast<LLVMFunctionType>(
-      cast<LLVMPointerType>(callee.getType()).getElementType());
-  SmallVector<Value> operands;
-  operands.reserve(1 + args.size());
-  operands.push_back(callee);
-  llvm::append_range(operands, args);
-  return build(builder, state, getCallOpResultTypes(calleeType),
-               TypeAttr::get(calleeType), FlatSymbolRefAttr(), operands,
-               /*fastmathFlags=*/nullptr, /*branch_weights=*/nullptr,
-               /*access_groups=*/nullptr, /*alias_scopes=*/nullptr,
-               /*noalias_scopes=*/nullptr, /*tbaa=*/nullptr);
-}
-
 CallInterfaceCallable CallOp::getCallableForCallee() {
   // Direct call.
   if (FlatSymbolRefAttr calleeAttr = getCalleeAttr())
@@ -1145,10 +978,7 @@ LogicalResult CallOp::verifySymbolUses(SymbolTableCollection &symbolTable) {
       return emitOpError("indirect call expects a pointer as callee: ")
              << getOperand(0).getType();
 
-    if (ptrType.isOpaque())
-      return success();
-
-    fnType = ptrType.getElementType();
+    return success();
   } else {
     Operation *callee =
         symbolTable.lookupNearestSymbolFrom(*this, calleeName.getAttr());
@@ -1848,17 +1678,6 @@ AddressOfOp::verifySymbolUses(SymbolTableCollection &symbolTable) {
     return emitOpError("pointer address space must match address space of the "
                        "referenced global");
 
-  if (type.isOpaque())
-    return success();
-
-  if (global && type.getElementType() != global.getType())
-    return emitOpError(
-        "the type must be a pointer to the type of the referenced global");
-
-  if (function && type.getElementType() != function.getFunctionType())
-    return emitOpError(
-        "the type must be a pointer to the type of the referenced function");
-
   return success();
 }
 
@@ -2135,9 +1954,6 @@ static bool isZeroAttribute(Attribute value) {
 }
 
 LogicalResult GlobalOp::verify() {
-  if (!LLVMPointerType::isValidElementType(getType()))
-    return emitOpError(
-        "expects type to be a valid element type for an LLVM pointer");
   if ((*this)->getParentOp() && !satisfiesLLVMModule((*this)->getParentOp()))
     return emitOpError("must appear at the module level");
 
@@ -2733,11 +2549,7 @@ void AtomicRMWOp::build(OpBuilder &builder, OperationState &state,
 }
 
 LogicalResult AtomicRMWOp::verify() {
-  auto ptrType = llvm::cast<LLVM::LLVMPointerType>(getPtr().getType());
   auto valType = getVal().getType();
-  if (!ptrType.isOpaque() && valType != ptrType.getElementType())
-    return emitOpError("expected LLVM IR element type for operand #0 to "
-                       "match type for operand #1");
   if (getBinOp() == AtomicBinOp::fadd || getBinOp() == AtomicBinOp::fsub ||
       getBinOp() == AtomicBinOp::fmin || getBinOp() == AtomicBinOp::fmax) {
     if (!mlir::LLVM::isCompatibleFloatingPointType(valType))
@@ -2790,9 +2602,6 @@ LogicalResult AtomicCmpXchgOp::verify() {
   if (!ptrType)
     return emitOpError("expected LLVM IR pointer type for operand #0");
   auto valType = getVal().getType();
-  if (!ptrType.isOpaque() && valType != ptrType.getElementType())
-    return emitOpError("expected LLVM IR element type for operand #0 to "
-                       "match type for all other operands");
   if (!isTypeCompatibleWithAtomicOp(valType,
                                     /*isPointerTypeAllowed=*/true))
     return emitOpError("unexpected LLVM IR type");
@@ -3185,14 +2994,7 @@ LogicalResult LLVMDialect::verifyParameterAttribute(Operation *op,
   auto checkPointerTypeMatches = [&]() -> LogicalResult {
     if (failed(checkPointerType()))
       return failure();
-    auto ptrType = llvm::cast<LLVMPointerType>(paramType);
-    auto typeAttr = llvm::cast<TypeAttr>(paramAttr.getValue());
 
-    if (!ptrType.isOpaque() && ptrType.getElementType() != typeAttr.getValue())
-      return op->emitError()
-             << name
-             << " attribute attached to LLVM pointer argument of "
-                "different type";
     return success();
   };
 
@@ -3317,8 +3119,7 @@ Operation *LLVMDialect::materializeConstant(OpBuilder &builder, Attribute value,
 
 Value mlir::LLVM::createGlobalString(Location loc, OpBuilder &builder,
                                      StringRef name, StringRef value,
-                                     LLVM::Linkage linkage,
-                                     bool useOpaquePointers) {
+                                     LLVM::Linkage linkage) {
   assert(builder.getInsertionBlock() &&
          builder.getInsertionBlock()->getParentOp() &&
          "expected builder to point to a block constrained in an op");
@@ -3334,19 +3135,11 @@ Value mlir::LLVM::createGlobalString(Location loc, OpBuilder &builder,
       loc, type, /*isConstant=*/true, linkage, name,
       builder.getStringAttr(value), /*alignment=*/0);
 
-  LLVMPointerType resultType;
-  LLVMPointerType charPtr;
-  if (!useOpaquePointers) {
-    resultType = LLVMPointerType::get(type);
-    charPtr = LLVMPointerType::get(IntegerType::get(ctx, 8));
-  } else {
-    resultType = charPtr = LLVMPointerType::get(ctx);
-  }
-
+  LLVMPointerType ptrType = LLVMPointerType::get(ctx);
   // Get the pointer to the first character in the global string.
-  Value globalPtr = builder.create<LLVM::AddressOfOp>(loc, resultType,
-                                                      global.getSymNameAttr());
-  return builder.create<LLVM::GEPOp>(loc, charPtr, type, globalPtr,
+  Value globalPtr =
+      builder.create<LLVM::AddressOfOp>(loc, ptrType, global.getSymNameAttr());
+  return builder.create<LLVM::GEPOp>(loc, ptrType, type, globalPtr,
                                      ArrayRef<GEPArg>{0, 0});
 }
 
diff --git a/mlir/lib/Dialect/LLVMIR/IR/LLVMMemorySlot.cpp b/mlir/lib/Dialect/LLVMIR/IR/LLVMMemorySlot.cpp
index 2cdc5d9fff84e78..cc2a917e58f6394 100644
--- a/mlir/lib/Dialect/LLVMIR/IR/LLVMMemorySlot.cpp
+++ b/mlir/lib/Dialect/LLVMIR/IR/LLVMMemorySlot.cpp
@@ -32,7 +32,7 @@ llvm::SmallVector<MemorySlot> LLVM::AllocaOp::getPromotableSlots() {
   if (!getOperation()->getBlock()->isEntryBlock())
     return {};
 
-  return {MemorySlot{getResult(), getResultPtrElementType()}};
+  return {MemorySlot{getResult(), getElemType()}};
 }
 
 Value LLVM::AllocaOp::getDefaultValue(const MemorySlot &slot,
@@ -61,8 +61,7 @@ SmallVector<DestructurableMemorySlot> LLVM::AllocaOp::getDestructurableSlots() {
   if (!mlir::matchPattern(getArraySize(), m_One()))
     return {};
 
-  Type elemType = getResultPtrElementType();
-  auto destructurable = dyn_cast<DestructurableTypeInterface>(elemType);
+  auto destructurable = dyn_cast<DestructurableTypeInterface>(getElemType());
   if (!destructurable)
     return {};
 
@@ -75,7 +74,8 @@ SmallVector<DestructurableMemorySlot> LLVM::AllocaOp::getDestructurableSlots() {
   for (Attribute index : llvm::make_first_range(destructuredType.value()))
     allocaTypeMap.insert({index, LLVM::LLVMPointerType::get(getContext())});
 
-  return {DestructurableMemorySlot{{getResult(), elemType}, {allocaTypeMap}}};
+  return {
+      DestructurableMemorySlot{{getResult(), getElemType()}, {allocaTypeMap}}};
 }
 
 DenseMap<Attribute, MemorySlot>
@@ -83,12 +83,9 @@ LLVM::AllocaOp::destructure(const DestructurableMemorySlot &slot,
                             const SmallPtrSetImpl<Attribute> &usedIndices,
                             RewriterBase &rewriter) {
   assert(slot.ptr == getResult());
-  Type elemType =
-      getElemType() ? *getElemType() : getResult().getType().getElementType();
-
   rewriter.setInsertionPointAfter(*this);
 
-  auto destructurableType = cast<DestructurableTypeInterface>(elemType);
+  auto destructurableType = cast<DestructurableTypeInterface>(getElemType());
   DenseMap<Attribute, MemorySlot> slotMap;
   for (Attribute index : usedIndices) {
     Type elemType = destructurableType.getTypeAtIndex(index);
@@ -337,11 +334,6 @@ bool LLVM::GEPOp::canRewire(const DestructurableMemorySlot &slot,
   if (!basePtrType)
     return false;
 
-  // Typed pointers are not supported. This should be removed once typed
-  // pointers are removed from the LLVM dialect.
-  if (!basePtrType.isOpaque())
-    return false;
-
   if (getBase() != slot.ptr || slot.elemType != getElemType())
     return false;
   if (!isFirstIndexZero(*this))
diff --git a/mlir/lib/Dialect/LLVMIR/IR/LLVMTypes.cpp b/mlir/lib/Dialect/LLVMIR/IR/LLVMTypes.cpp
index bc8300a8b7329ea..8841aa8362569a0 100644
--- a/mlir/lib/Dialect/LLVMIR/IR/LLVMTypes.cpp
+++ b/mlir/lib/Dialect/LLVMIR/IR/LLVMTypes.cpp
@@ -75,40 +75,6 @@ static void printFunctionTypes(AsmPrinter &p, ArrayRef<Type> params,
   p << ')';
 }
 
-//===----------------------------------------------------------------------===//
-// custom<Pointer>
-//===----------------------------------------------------------------------===//
-
-static ParseResult parsePointer(AsmParser &p, Type &elementType,
-                                unsigned &addressSpace) {
-  // `<` 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;
-  }
-}
-
 //===----------------------------------------------------------------------===//
 // custom<ExtTypeParams>
 //===----------------------------------------------------------------------===//
@@ -285,33 +251,6 @@ LLVMFunctionType::verify(function_ref<InFlightDiagnostic()> emitError,
   return success();
 }
 
-//===----------------------------------------------------------------------===//
-// LLVMPointerType
-//===----------------------------------------------------------------------===//
-
-bool LLVMPointerType::isValidElementType(Type type) {
-  if (!type)
-    return true;
-  return isCompatibleOuterType(type)
-             ? !llvm::isa<LLVMVoidType, LLVMTokenType, LLVMMetadataType,
-                          LLVMLabelType>(type)
-             : llvm::isa<PointerElementTypeInterface>(type);
-}
-
-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);
-}
-
-LogicalResult
-LLVMPointerType::verify(function_ref<InFlightDiagnostic()> emitError,
-                        Type pointee, unsigned) {
-  if (!isValidElementType(pointee))
-    return emitError() << "invalid pointer element type: " << pointee;
-  return success();
-}
-
 //===----------------------------------------------------------------------===//
 // DataLayoutTypeInterface
 
@@ -369,9 +308,7 @@ LLVMPointerType::getTypeSizeInBits(const DataLayout &dataLayout,
 
   // For other memory spaces, use the size of the pointer to the default memory
   // space.
-  if (isOpaque())
-    return dataLayout.getTypeSizeInBits(get(getContext()));
-  return dataLayout.getTypeSizeInBits(get(getElementType()));
+  return dataLayout.getTypeSizeInBits(get(getContext()));
 }
 
 unsigned LLVMPointerType::getABIAlignment(const DataLayout &dataLayout,
@@ -380,9 +317,7 @@ unsigned LLVMPointerType::getABIAlignment(const DataLayout &dataLayout,
           getPointerDataLayoutEntry(params, *this, PtrDLEntryPos::Abi))
     return *alignment;
 
-  if (isOpaque())
-    return dataLayout.getTypeABIAlignment(get(getContext()));
-  return dataLayout.getTypeABIAlignment(get(getElementType()));
+  return dataLayout.getTypeABIAlignment(get(getContext()));
 }
 
 unsigned
@@ -392,9 +327,7 @@ LLVMPointerType::getPreferredAlignment(const DataLayout &dataLayout,
           getPointerDataLayoutEntry(params, *this, PtrDLEntryPos::Preferred))
     return *alignment;
 
-  if (isOpaque())
-    return dataLayout.getTypePreferredAlignment(get(getContext()));
-  return dataLayout.getTypePreferredAlignment(get(getElementType()));
+  return dataLayout.getTypePreferredAlignment(get(getContext()));
 }
 
 bool LLVMPointerType::areCompatible(DataLayoutEntryListRef oldLayout,
@@ -440,7 +373,6 @@ LogicalResult LLVMPointerType::verifyEntries(DataLayoutEntryListRef entries,
   for (DataLayoutEntryInterface entry : entries) {
     if (!entry.isTypeEntry())
       continue;
-    auto key = llvm::cast<LLVMPointerType>(entry.getKey().get<Type>());
     auto values = llvm::dyn_cast<DenseIntElementsAttr>(entry.getValue());
     if (!values || (values.size() != 3 && values.size() != 4)) {
       return emitError(loc)
@@ -448,10 +380,6 @@ LogicalResult LLVMPointerType::verifyEntries(DataLayoutEntryListRef entries,
              << " to be a dense integer elements attribute with 3 or 4 "
                 "elements";
     }
-    if (key.getElementType() && !key.getElementType().isInteger(8)) {
-      return emitError(loc) << "unexpected layout attribute for pointer to "
-                            << key.getElementType();
-    }
     if (extractPointerSpecValue(values, PtrDLEntryPos::Abi) >
         extractPointerSpecValue(values, PtrDLEntryPos::Preferred)) {
       return emitError(loc) << "preferred alignment is expected to be at least "
@@ -869,11 +797,7 @@ static bool isCompatibleImpl(Type type, DenseSet<Type> &compatibleTypes) {
             return vecType.getRank() == 1 &&
                    isCompatible(vecType.getElementType());
           })
-          .Case<LLVMPointerType>([&](auto pointerType) {
-            if (pointerType.isOpaque())
-              return true;
-            return isCompatible(pointerType.getElementType());
-          })
+          .Case<LLVMPointerType>([&](auto pointerType) { return true; })
           .Case<LLVMTargetExtType>([&](auto extType) {
             return llvm::all_of(extType.getTypeParams(), isCompatible);
           })
diff --git a/mlir/lib/Dialect/LLVMIR/Transforms/TypeConsistency.cpp b/mlir/lib/Dialect/LLVMIR/Transforms/TypeConsistency.cpp
index 9731689e551762b..ee491a6c558b888 100644
--- a/mlir/lib/Dialect/LLVMIR/Transforms/TypeConsistency.cpp
+++ b/mlir/lib/Dialect/LLVMIR/Transforms/TypeConsistency.cpp
@@ -92,10 +92,6 @@ LogicalResult AddFieldGetterToStructDirectUse<LoadOp>::matchAndRewrite(
     LoadOp load, PatternRewriter &rewriter) const {
   PatternRewriter::InsertionGuard guard(rewriter);
 
-  // Load from typed pointers are not supported.
-  if (!load.getAddr().getType().isOpaque())
-    return failure();
-
   Type inconsistentElementType =
       isElementTypeInconsistent(load.getAddr(), load.getType());
   if (!inconsistentElementType)
@@ -129,10 +125,6 @@ LogicalResult AddFieldGetterToStructDirectUse<StoreOp>::matchAndRewrite(
     StoreOp store, PatternRewriter &rewriter) const {
   PatternRewriter::InsertionGuard guard(rewriter);
 
-  // Store to typed pointers are not supported.
-  if (!store.getAddr().getType().isOpaque())
-    return failure();
-
   Type inconsistentElementType =
       isElementTypeInconsistent(store.getAddr(), store.getValue().getType());
   if (!inconsistentElementType)
diff --git a/mlir/lib/Dialect/OpenACC/IR/OpenACC.cpp b/mlir/lib/Dialect/OpenACC/IR/OpenACC.cpp
index df64d561f46cb3e..ff2eb9ebfc42f5e 100644
--- a/mlir/lib/Dialect/OpenACC/IR/OpenACC.cpp
+++ b/mlir/lib/Dialect/OpenACC/IR/OpenACC.cpp
@@ -27,14 +27,19 @@ using namespace acc;
 #include "mlir/Dialect/OpenACC/OpenACCTypeInterfaces.cpp.inc"
 
 namespace {
-/// Model for pointer-like types that already provide a `getElementType` method.
-template <typename T>
-struct PointerLikeModel
-    : public PointerLikeType::ExternalModel<PointerLikeModel<T>, T> {
+struct MemRefPointerLikeModel
+    : public PointerLikeType::ExternalModel<MemRefPointerLikeModel,
+                                            MemRefType> {
   Type getElementType(Type pointer) const {
-    return llvm::cast<T>(pointer).getElementType();
+    return llvm::cast<MemRefType>(pointer).getElementType();
   }
 };
+
+struct LLVMPointerPointerLikeModel
+    : public PointerLikeType::ExternalModel<LLVMPointerPointerLikeModel,
+                                            LLVM::LLVMPointerType> {
+  Type getElementType(Type pointer) const { return Type(); }
+};
 } // namespace
 
 //===----------------------------------------------------------------------===//
@@ -58,9 +63,9 @@ void OpenACCDialect::initialize() {
   // By attaching interfaces here, we make the OpenACC dialect dependent on
   // the other dialects. This is probably better than having dialects like LLVM
   // and memref be dependent on OpenACC.
-  LLVM::LLVMPointerType::attachInterface<
-      PointerLikeModel<LLVM::LLVMPointerType>>(*getContext());
-  MemRefType::attachInterface<PointerLikeModel<MemRefType>>(*getContext());
+  MemRefType::attachInterface<MemRefPointerLikeModel>(*getContext());
+  LLVM::LLVMPointerType::attachInterface<LLVMPointerPointerLikeModel>(
+      *getContext());
 }
 
 //===----------------------------------------------------------------------===//
@@ -1023,17 +1028,13 @@ void EnterDataOp::getCanonicalizationPatterns(RewritePatternSet &results,
 // AtomicReadOp
 //===----------------------------------------------------------------------===//
 
-LogicalResult AtomicReadOp::verify() {
-  return verifyCommon();
-}
+LogicalResult AtomicReadOp::verify() { return verifyCommon(); }
 
 //===----------------------------------------------------------------------===//
 // AtomicWriteOp
 //===----------------------------------------------------------------------===//
 
-LogicalResult AtomicWriteOp::verify() {
-  return verifyCommon();
-}
+LogicalResult AtomicWriteOp::verify() { return verifyCommon(); }
 
 //===----------------------------------------------------------------------===//
 // AtomicUpdateOp
@@ -1054,13 +1055,9 @@ LogicalResult AtomicUpdateOp::canonicalize(AtomicUpdateOp op,
   return failure();
 }
 
-LogicalResult AtomicUpdateOp::verify() {
-  return verifyCommon();
-}
+LogicalResult AtomicUpdateOp::verify() { return verifyCommon(); }
 
-LogicalResult AtomicUpdateOp::verifyRegions() {
-  return verifyRegionsCommon();
-}
+LogicalResult AtomicUpdateOp::verifyRegions() { return verifyRegionsCommon(); }
 
 //===----------------------------------------------------------------------===//
 // AtomicCaptureOp
@@ -1084,9 +1081,7 @@ AtomicUpdateOp AtomicCaptureOp::getAtomicUpdateOp() {
   return dyn_cast<AtomicUpdateOp>(getSecondOp());
 }
 
-LogicalResult AtomicCaptureOp::verifyRegions() {
-  return verifyRegionsCommon();
-}
+LogicalResult AtomicCaptureOp::verifyRegions() { return verifyRegionsCommon(); }
 
 //===----------------------------------------------------------------------===//
 // DeclareEnterOp
diff --git a/mlir/lib/Dialect/OpenMP/IR/OpenMPDialect.cpp b/mlir/lib/Dialect/OpenMP/IR/OpenMPDialect.cpp
index 3b792a26d1823fe..f6757aba664f985 100644
--- a/mlir/lib/Dialect/OpenMP/IR/OpenMPDialect.cpp
+++ b/mlir/lib/Dialect/OpenMP/IR/OpenMPDialect.cpp
@@ -39,15 +39,20 @@ using namespace mlir;
 using namespace mlir::omp;
 
 namespace {
-/// Model for pointer-like types that already provide a `getElementType` method.
-template <typename T>
-struct PointerLikeModel
-    : public PointerLikeType::ExternalModel<PointerLikeModel<T>, T> {
+struct MemRefPointerLikeModel
+    : public PointerLikeType::ExternalModel<MemRefPointerLikeModel,
+                                            MemRefType> {
   Type getElementType(Type pointer) const {
-    return llvm::cast<T>(pointer).getElementType();
+    return llvm::cast<MemRefType>(pointer).getElementType();
   }
 };
 
+struct LLVMPointerPointerLikeModel
+    : public PointerLikeType::ExternalModel<LLVMPointerPointerLikeModel,
+                                            LLVM::LLVMPointerType> {
+  Type getElementType(Type pointer) const { return Type(); }
+};
+
 struct OpenMPDialectFoldInterface : public DialectFoldInterface {
   using DialectFoldInterface::DialectFoldInterface;
 
@@ -73,11 +78,9 @@ void OpenMPDialect::initialize() {
       >();
 
   addInterface<OpenMPDialectFoldInterface>();
-  LLVM::LLVMPointerType::attachInterface<
-      PointerLikeModel<LLVM::LLVMPointerType>>(*getContext());
-  MemRefType::attachInterface<PointerLikeModel<MemRefType>>(*getContext());
-  LLVM::LLVMPointerType::attachInterface<
-      PointerLikeModel<LLVM::LLVMPointerType>>(*getContext());
+  MemRefType::attachInterface<MemRefPointerLikeModel>(*getContext());
+  LLVM::LLVMPointerType::attachInterface<LLVMPointerPointerLikeModel>(
+      *getContext());
 
   // Attach default offload module interface to module op to access
   // offload functionality through
@@ -1342,9 +1345,7 @@ LogicalResult AtomicUpdateOp::verify() {
   return verifySynchronizationHint(*this, getHintVal());
 }
 
-LogicalResult AtomicUpdateOp::verifyRegions() {
-  return verifyRegionsCommon();
-}
+LogicalResult AtomicUpdateOp::verifyRegions() { return verifyRegionsCommon(); }
 
 //===----------------------------------------------------------------------===//
 // Verifier for AtomicCaptureOp
diff --git a/mlir/test/Dialect/LLVMIR/global.mlir b/mlir/test/Dialect/LLVMIR/global.mlir
index a33fff3967e4d59..e5d7d6d8f8c40f8 100644
--- a/mlir/test/Dialect/LLVMIR/global.mlir
+++ b/mlir/test/Dialect/LLVMIR/global.mlir
@@ -117,11 +117,6 @@ llvm.mlir.global internal protected unnamed_addr @protected(42 : i32) : i32
 
 // -----
 
-// expected-error @+1 {{expects type to be a valid element type for an LLVM pointer}}
-llvm.mlir.global internal constant @constant(37.0) : !llvm.label
-
-// -----
-
 // expected-error @+1 {{'addr_space' failed to satisfy constraint: 32-bit signless integer attribute whose value is non-negative}}
 "llvm.mlir.global"() ({}) {sym_name = "foo", global_type = i64, value = 42 : i64, addr_space = -1 : i32, linkage = #llvm.linkage<private>} : () -> ()
 
diff --git a/mlir/test/Dialect/LLVMIR/invalid.mlir b/mlir/test/Dialect/LLVMIR/invalid.mlir
index fe2f94454561a08..1d51796abb03faf 100644
--- a/mlir/test/Dialect/LLVMIR/invalid.mlir
+++ b/mlir/test/Dialect/LLVMIR/invalid.mlir
@@ -90,30 +90,23 @@ func.func @alloca_non_integer_alignment() {
 
 // -----
 
-func.func @alloca_opaque_ptr_no_type(%sz : i64) {
-  // expected-error at below {{expected 'elem_type' attribute if opaque pointer type is used}}
-  "llvm.alloca"(%sz) : (i64) -> !llvm.ptr
-}
-
-// -----
-
 func.func @gep_missing_input_result_type(%pos : i64, %base : !llvm.ptr) {
   // expected-error at +1 {{2 operands present, but expected 0}}
-  llvm.getelementptr %base[%pos] : () -> ()
+  llvm.getelementptr %base[%pos] : () -> (), i64
 }
 
 // -----
 
 func.func @gep_missing_input_type(%pos : i64, %base : !llvm.ptr) {
   // expected-error at +1 {{2 operands present, but expected 0}}
-  llvm.getelementptr %base[%pos] : () -> (!llvm.ptr)
+  llvm.getelementptr %base[%pos] : () -> (!llvm.ptr), i64
 }
 
 // -----
 
 func.func @gep_missing_result_type(%pos : i64, %base : !llvm.ptr) {
   // expected-error at +1 {{op requires one result}}
-  llvm.getelementptr %base[%pos] : (!llvm.ptr, i64) -> ()
+  llvm.getelementptr %base[%pos] : (!llvm.ptr, i64) -> (), i64
 }
 
 // -----
@@ -132,20 +125,6 @@ func.func @gep_too_few_dynamic(%base : !llvm.ptr) {
 
 // -----
 
-func.func @load_non_llvm_type(%foo : memref<f32>) {
-  // expected-error at +1 {{expected LLVM pointer type}}
-  llvm.load %foo : memref<f32>
-}
-
-// -----
-
-func.func @load_non_ptr_type(%foo : f32) {
-  // expected-error at +1 {{expected LLVM pointer type}}
-  llvm.load %foo : f32
-}
-
-// -----
-
 func.func @load_syncscope(%ptr : !llvm.ptr) {
   // expected-error at below {{expected syncscope to be null for non-atomic access}}
   %1 = "llvm.load"(%ptr) {syncscope = "singlethread"} : (!llvm.ptr) -> (f32)
@@ -181,27 +160,6 @@ func.func @load_unaligned_atomic(%ptr : !llvm.ptr) {
 
 // -----
 
-func.func @store_non_llvm_type(%foo : memref<f32>, %bar : f32) {
-  // expected-error at +1 {{expected LLVM pointer type}}
-  llvm.store %bar, %foo : memref<f32>
-}
-
-// -----
-
-func.func @store_non_ptr_type(%foo : f32, %bar : f32) {
-  // expected-error at +1 {{expected LLVM pointer type}}
-  llvm.store %bar, %foo : f32
-}
-
-// -----
-
-func.func @store_malformed_elem_type(%foo: !llvm.ptr, %bar: f32) {
-  // expected-error at +1 {{expected non-function type}}
-  llvm.store %bar, %foo : !llvm.ptr, "f32"
-}
-
-// -----
-
 func.func @store_syncscope(%val : f32, %ptr : !llvm.ptr) {
   // expected-error at below {{expected syncscope to be null for non-atomic access}}
   "llvm.store"(%val, %ptr) {syncscope = "singlethread"} : (f32, !llvm.ptr) -> ()
@@ -632,14 +590,6 @@ func.func @nvvm_invalid_mma_8(%a0 : i32, %a1 : i32,
 
 // -----
 
-func.func @atomicrmw_expected_ptr(%f32 : f32) {
-  // expected-error at +1 {{operand #0 must be LLVM pointer to floating point LLVM type or LLVM pointer 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 : f32) {
   // expected-error at +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) -> i32
@@ -672,14 +622,6 @@ func.func @atomicrmw_expected_int(%f32_ptr : !llvm.ptr, %f32 : f32) {
 
 // -----
 
-func.func @cmpxchg_expected_ptr(%f32 : f32) {
-  // expected-error at +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_value_operands(%ptr : !llvm.ptr, %i32 : i32, %i64 : i64) {
   // expected-error at +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)>
diff --git a/mlir/unittests/Dialect/LLVMIR/LLVMTypeTest.cpp b/mlir/unittests/Dialect/LLVMIR/LLVMTypeTest.cpp
index aa19b5c651f5594..083dec819a0e054 100644
--- a/mlir/unittests/Dialect/LLVMIR/LLVMTypeTest.cpp
+++ b/mlir/unittests/Dialect/LLVMIR/LLVMTypeTest.cpp
@@ -17,37 +17,3 @@ TEST_F(LLVMIRTest, IsStructTypeMutable) {
   ASSERT_TRUE(bool(structTy));
   ASSERT_TRUE(structTy.hasTrait<TypeTrait::IsMutable>());
 }
-
-TEST_F(LLVMIRTest, MutualReferencedSubElementTypes) {
-  auto fooStructTy = LLVMStructType::getIdentified(&context, "foo");
-  ASSERT_TRUE(bool(fooStructTy));
-  auto barStructTy = LLVMStructType::getIdentified(&context, "bar");
-  ASSERT_TRUE(bool(barStructTy));
-
-  // Created two structs that are referencing each other.
-  Type fooBody[] = {LLVMPointerType::get(barStructTy)};
-  ASSERT_TRUE(succeeded(fooStructTy.setBody(fooBody, /*isPacked=*/false)));
-  Type barBody[] = {LLVMPointerType::get(fooStructTy)};
-  ASSERT_TRUE(succeeded(barStructTy.setBody(barBody, /*isPacked=*/false)));
-
-  // Test if walkSubElements goes into infinite loops.
-  SmallVector<Type, 4> subElementTypes;
-  fooStructTy.walk([&](Type type) { subElementTypes.push_back(type); });
-  ASSERT_EQ(subElementTypes.size(), 4U);
-
-  // !llvm.ptr<struct<"foo",...>>
-  ASSERT_TRUE(isa<LLVMPointerType>(subElementTypes[0]));
-
-  // !llvm.struct<"bar",...>
-  auto structType = dyn_cast<LLVMStructType>(subElementTypes[1]);
-  ASSERT_TRUE(bool(structType));
-  ASSERT_TRUE(structType.getName().equals("bar"));
-
-  // !llvm.ptr<struct<"bar",...>>
-  ASSERT_TRUE(isa<LLVMPointerType>(subElementTypes[2]));
-
-  // !llvm.struct<"foo",...>
-  structType = dyn_cast<LLVMStructType>(subElementTypes[3]);
-  ASSERT_TRUE(bool(structType));
-  ASSERT_TRUE(structType.getName().equals("foo"));
-}

>From f78ccce6a2ab16b21131a442d3a9f19926e446e9 Mon Sep 17 00:00:00 2001
From: Christian Ulmann <christianulmann at gmail.com>
Date: Sat, 4 Nov 2023 19:46:02 +0100
Subject: [PATCH 2/2] address review comments

---
 mlir/include/mlir/Dialect/LLVMIR/LLVMOpBase.td | 9 +++++----
 mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td    | 4 ++--
 mlir/include/mlir/Dialect/LLVMIR/ROCDLOps.td   | 2 +-
 mlir/lib/Dialect/LLVMIR/IR/LLVMDialect.cpp     | 7 +++++++
 mlir/test/Dialect/LLVMIR/global.mlir           | 4 ++++
 mlir/test/Dialect/LLVMIR/invalid.mlir          | 7 +++++++
 6 files changed, 26 insertions(+), 7 deletions(-)

diff --git a/mlir/include/mlir/Dialect/LLVMIR/LLVMOpBase.td b/mlir/include/mlir/Dialect/LLVMIR/LLVMOpBase.td
index 0166fbb647b41b5..a459000581028ae 100644
--- a/mlir/include/mlir/Dialect/LLVMIR/LLVMOpBase.td
+++ b/mlir/include/mlir/Dialect/LLVMIR/LLVMOpBase.td
@@ -56,10 +56,11 @@ def LLVM_AnyPointer : Type<CPred<"::llvm::isa<::mlir::LLVM::LLVMPointerType>($_s
                           "LLVM pointer type", "::mlir::LLVM::LLVMPointerType">;
 
 // Opaque pointer in a given address space.
-class LLVM_OpaquePointerInAddressSpace<int addressSpace> : Type<
-        CPred<
-          "::llvm::cast<::mlir::LLVM::LLVMPointerType>($_self).getAddressSpace() == "
-          # addressSpace>,
+class LLVM_PointerInAddressSpace<int addressSpace> : Type<
+  And<[LLVM_AnyPointer.predicate,
+      CPred<
+        "::llvm::cast<::mlir::LLVM::LLVMPointerType>($_self).getAddressSpace() == "
+        # addressSpace>]>,
   "Opaque LLVM pointer in address space " # addressSpace,
   "::mlir::LLVM::LLVMPointerType"> {
   let builderCall = "$_builder.getType<::mlir::LLVM::LLVMPointerType>("
diff --git a/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td b/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
index 406519d3eaa7272..16f66891a050845 100644
--- a/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
+++ b/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
@@ -19,8 +19,8 @@ include "mlir/Dialect/LLVMIR/LLVMOpBase.td"
 include "mlir/Interfaces/SideEffectInterfaces.td"
 include "mlir/Dialect/LLVMIR/BasicPtxBuilderInterface.td"
 
-def LLVM_ptr_global : LLVM_OpaquePointerInAddressSpace<1>;
-def LLVM_ptr_shared : LLVM_OpaquePointerInAddressSpace<3>;
+def LLVM_ptr_global : LLVM_PointerInAddressSpace<1>;
+def LLVM_ptr_shared : LLVM_PointerInAddressSpace<3>;
 
 //===----------------------------------------------------------------------===//
 // NVVM dialect definitions
diff --git a/mlir/include/mlir/Dialect/LLVMIR/ROCDLOps.td b/mlir/include/mlir/Dialect/LLVMIR/ROCDLOps.td
index 6c6419bf238b457..48b830ae34f2922 100644
--- a/mlir/include/mlir/Dialect/LLVMIR/ROCDLOps.td
+++ b/mlir/include/mlir/Dialect/LLVMIR/ROCDLOps.td
@@ -275,7 +275,7 @@ def ROCDL_wmma_i32_16x16x16_iu4 : ROCDL_Wmma_IntrOp<"wmma.i32.16x16x16.iu4">;
 // raw buffer mode).
 //===---------------------------------------------------------------------===//
 
-def ROCDLBufferRsrc : LLVM_OpaquePointerInAddressSpace<8>;
+def ROCDLBufferRsrc : LLVM_PointerInAddressSpace<8>;
 
 def ROCDL_MakeBufferRsrcOp :
   ROCDL_IntrOp<"make.buffer.rsrc", [], [0], [Pure], 1>,
diff --git a/mlir/lib/Dialect/LLVMIR/IR/LLVMDialect.cpp b/mlir/lib/Dialect/LLVMIR/IR/LLVMDialect.cpp
index 3ce7e9abccaa387..33af73d8a4ac5fb 100644
--- a/mlir/lib/Dialect/LLVMIR/IR/LLVMDialect.cpp
+++ b/mlir/lib/Dialect/LLVMIR/IR/LLVMDialect.cpp
@@ -1954,6 +1954,13 @@ static bool isZeroAttribute(Attribute value) {
 }
 
 LogicalResult GlobalOp::verify() {
+  bool validType = isCompatibleOuterType(getType())
+                       ? !llvm::isa<LLVMVoidType, LLVMTokenType,
+                                    LLVMMetadataType, LLVMLabelType>(getType())
+                       : llvm::isa<PointerElementTypeInterface>(getType());
+  if (!validType)
+    return emitOpError(
+        "expects type to be a valid element type for an LLVM global");
   if ((*this)->getParentOp() && !satisfiesLLVMModule((*this)->getParentOp()))
     return emitOpError("must appear at the module level");
 
diff --git a/mlir/test/Dialect/LLVMIR/global.mlir b/mlir/test/Dialect/LLVMIR/global.mlir
index e5d7d6d8f8c40f8..81178b2ef901f4f 100644
--- a/mlir/test/Dialect/LLVMIR/global.mlir
+++ b/mlir/test/Dialect/LLVMIR/global.mlir
@@ -117,6 +117,10 @@ llvm.mlir.global internal protected unnamed_addr @protected(42 : i32) : i32
 
 // -----
 
+// expected-error @+1 {{expects type to be a valid element type for an LLVM global}}
+llvm.mlir.global internal constant @constant(37.0) : !llvm.label
+
+// -----
 // expected-error @+1 {{'addr_space' failed to satisfy constraint: 32-bit signless integer attribute whose value is non-negative}}
 "llvm.mlir.global"() ({}) {sym_name = "foo", global_type = i64, value = 42 : i64, addr_space = -1 : i32, linkage = #llvm.linkage<private>} : () -> ()
 
diff --git a/mlir/test/Dialect/LLVMIR/invalid.mlir b/mlir/test/Dialect/LLVMIR/invalid.mlir
index 1d51796abb03faf..d6960f00f0428cb 100644
--- a/mlir/test/Dialect/LLVMIR/invalid.mlir
+++ b/mlir/test/Dialect/LLVMIR/invalid.mlir
@@ -125,6 +125,13 @@ func.func @gep_too_few_dynamic(%base : !llvm.ptr) {
 
 // -----
 
+func.func @load_non_llvm_type(%foo : memref<f32>) {
+  // expected-error at +1 {{op operand #0 must be LLVM pointer type}}
+  llvm.load %foo : memref<f32> -> f32
+}
+
+// -----
+
 func.func @load_syncscope(%ptr : !llvm.ptr) {
   // expected-error at below {{expected syncscope to be null for non-atomic access}}
   %1 = "llvm.load"(%ptr) {syncscope = "singlethread"} : (!llvm.ptr) -> (f32)



More information about the Mlir-commits mailing list