diff --git a/llvm/include/llvm/IR/Verifier.h b/llvm/include/llvm/IR/Verifier.h
index 8dbb9c8a41d7e..135c6ab4ebb1f 100644
--- a/llvm/include/llvm/IR/Verifier.h
+++ b/llvm/include/llvm/IR/Verifier.h
@@ -21,20 +21,135 @@
 #define LLVM_IR_VERIFIER_H
 
 #include "llvm/ADT/DenseMap.h"
+#include "llvm/IR/DebugProgramInstruction.h"
+#include "llvm/IR/Metadata.h"
+#include "llvm/IR/ModuleSlotTracker.h"
 #include "llvm/IR/PassManager.h"
 #include "llvm/Support/Compiler.h"
+#include "llvm/Support/Printable.h"
 #include <utility>
 
 namespace llvm {
 
 class APInt;
+class Attribute;
+class AttributeList;
+class AttributeSet;
+class CallBase;
+class Comdat;
+class DataLayout;
 class Function;
 class FunctionPass;
 class Instruction;
-class MDNode;
+class LLVMContext;
 class Module;
+class Triple;
+class VerifierSupport;
 class raw_ostream;
-struct VerifierSupport;
+
+/// Base class for IR verifier plugins.
+///
+/// To add a plugin, derive from this class and then instantiate it once.
+class VerifierPlugin {
+public:
+  VerifierPlugin();
+  virtual ~VerifierPlugin();
+
+  /// Called when the verifier finds a call (or invoke) to an intrinsic it
+  /// doesn't understand.
+  ///
+  /// If the plugin recognizes the intrinsic, it should report any verifier
+  /// errors via the given helper object.
+  virtual void verifyIntrinsicCall(CallBase &Call, VerifierSupport &VS) const;
+};
+
+class VerifierSupport {
+public:
+  raw_ostream *OS;
+  const Module &M;
+  ModuleSlotTracker MST;
+  const Triple &TT;
+  const DataLayout &DL;
+  LLVMContext &Context;
+
+  /// Track the brokenness of the module while recursively visiting.
+  bool Broken = false;
+  /// Broken debug info can be "recovered" from by stripping the debug info.
+  bool BrokenDebugInfo = false;
+  /// Whether to treat broken debug info as an error.
+  bool TreatBrokenDebugInfoAsError = true;
+
+  explicit VerifierSupport(raw_ostream *OS, const Module &M);
+
+private:
+  LLVM_ABI void Write(const Module *M);
+  LLVM_ABI void Write(const Value *V);
+  LLVM_ABI void Write(const Value &V);
+  LLVM_ABI void Write(const DbgRecord *DR);
+  LLVM_ABI void Write(DbgVariableRecord::LocationType Type);
+  LLVM_ABI void Write(const Metadata *MD);
+
+  template <class T> void Write(const MDTupleTypedArrayWrapper<T> &MD) {
+    Write(MD.get());
+  }
+
+  LLVM_ABI void Write(const NamedMDNode *NMD);
+  LLVM_ABI void Write(Type *T);
+  LLVM_ABI void Write(const Comdat *C);
+  LLVM_ABI void Write(const APInt *AI);
+  LLVM_ABI void Write(const unsigned i) { *OS << i << '\n'; }
+
+  // NOLINTNEXTLINE(readability-identifier-naming)
+  LLVM_ABI void Write(const Attribute *A);
+  // NOLINTNEXTLINE(readability-identifier-naming)
+  LLVM_ABI void Write(const AttributeSet *AS);
+  // NOLINTNEXTLINE(readability-identifier-naming)
+  LLVM_ABI void Write(const AttributeList *AL);
+  LLVM_ABI void Write(Printable P) { *OS << P << '\n'; }
+
+  template <typename T> void Write(ArrayRef<T> Vs) {
+    for (const T &V : Vs)
+      Write(V);
+  }
+
+  template <typename T1, typename... Ts>
+  void WriteTs(const T1 &V1, const Ts &...Vs) {
+    Write(V1);
+    WriteTs(Vs...);
+  }
+
+  template <typename... Ts> void WriteTs() {}
+
+public:
+  /// A check failed, so printout out the condition and the message.
+  ///
+  /// This provides a nice place to put a breakpoint if you want to see why
+  /// something is not correct.
+  LLVM_ABI void CheckFailed(const Twine &Message);
+
+  /// A check failed (with values to print).
+  ///
+  /// This calls the Message-only version so that the above is easier to set a
+  /// breakpoint on.
+  template <typename T1, typename... Ts>
+  void CheckFailed(const Twine &Message, const T1 &V1, const Ts &...Vs) {
+    CheckFailed(Message);
+    if (OS)
+      WriteTs(V1, Vs...);
+  }
+
+  /// A debug info check failed.
+  LLVM_ABI void DebugInfoCheckFailed(const Twine &Message);
+
+  /// A debug info check failed (with values to print).
+  template <typename T1, typename... Ts>
+  void DebugInfoCheckFailed(const Twine &Message, const T1 &V1,
+                            const Ts &...Vs) {
+    DebugInfoCheckFailed(Message);
+    if (OS)
+      WriteTs(V1, Vs...);
+  }
+};
 
 /// Verify that the TBAA Metadatas are valid.
 class TBAAVerifier {
diff --git a/llvm/lib/IR/Verifier.cpp b/llvm/lib/IR/Verifier.cpp
index c06b60fd2d9a9..0b393f39e75ff 100644
--- a/llvm/lib/IR/Verifier.cpp
+++ b/llvm/lib/IR/Verifier.cpp
@@ -119,6 +119,7 @@
 #include "llvm/Support/ErrorHandling.h"
 #include "llvm/Support/MathExtras.h"
 #include "llvm/Support/ModRef.h"
+#include "llvm/Support/Mutex.h"
 #include "llvm/Support/TimeProfiler.h"
 #include "llvm/Support/raw_ostream.h"
 #include <algorithm>
@@ -127,6 +128,7 @@
 #include <memory>
 #include <optional>
 #include <string>
+#include <thread>
 #include <utility>
 
 using namespace llvm;
@@ -136,194 +138,307 @@ static cl::opt<bool> VerifyNoAliasScopeDomination(
     cl::desc("Ensure that llvm.experimental.noalias.scope.decl for identical "
              "scopes are not dominating"));
 
-namespace llvm {
+namespace {
 
-struct VerifierSupport {
-  raw_ostream *OS;
-  const Module &M;
-  ModuleSlotTracker MST;
-  const Triple &TT;
-  const DataLayout &DL;
-  LLVMContext &Context;
+class PluginRegistryLock;
+class PluginRegistryReader;
 
-  /// Track the brokenness of the module while recursively visiting.
-  bool Broken = false;
-  /// Broken debug info can be "recovered" from by stripping the debug info.
-  bool BrokenDebugInfo = false;
-  /// Whether to treat broken debug info as an error.
-  bool TreatBrokenDebugInfoAsError = true;
+/// Registry for verifier plugins.
+///
+/// The registry satifies the following implementation constraints:
+///
+///  * Support dynamically loading and unloading plugins from a thread (e.g.
+///    from dlopen()/dlclose()) while another thread may be in the verifier
+///  * Fast path for iterating over plugins that is lock-free and avoids
+///    cache-line ping pong
+///  * Plugin teardown may happen due to report_fatal_error from a thread that
+///    is currently in the verifier
+///
+/// The implementation achieves this by registering a "reader" object while the
+/// verifier is active. The reader object holds a hazard pointer to the plugins
+/// list that it is currently using.
+class PluginRegistry {
+  friend PluginRegistryReader;
+  friend PluginRegistryLock;
 
-  explicit VerifierSupport(raw_ostream *OS, const Module &M)
-      : OS(OS), M(M), MST(&M), TT(M.getTargetTriple()), DL(M.getDataLayout()),
-        Context(M.getContext()) {}
+  using List = SmallVector<const VerifierPlugin *>;
 
-private:
-  void Write(const Module *M) {
-    *OS << "; ModuleID = '" << M->getModuleIdentifier() << "'\n";
-  }
+  sys::Mutex Mutex;
+  SmallVector<PluginRegistryReader *> Readers;
+  List PluginsStorage[2];
+  std::atomic<List *> Plugins;
 
-  void Write(const Value *V) {
-    if (V)
-      Write(*V);
+  PluginRegistry() {
+    Plugins.store(&PluginsStorage[0], std::memory_order_relaxed);
   }
 
-  void Write(const Value &V) {
-    if (isa<Instruction>(V)) {
-      V.print(*OS, MST);
-      *OS << '\n';
-    } else {
-      V.printAsOperand(*OS, true, MST);
-      *OS << '\n';
-    }
-  }
+  template <typename FnT> void updatePlugins(FnT &&F);
 
-  void Write(const DbgRecord *DR) {
-    if (DR) {
-      DR->print(*OS, MST, false);
-      *OS << '\n';
-    }
+public:
+  static PluginRegistry &get() {
+    static PluginRegistry R;
+    return R;
   }
 
-  void Write(DbgVariableRecord::LocationType Type) {
-    switch (Type) {
-    case DbgVariableRecord::LocationType::Value:
-      *OS << "value";
-      break;
-    case DbgVariableRecord::LocationType::Declare:
-      *OS << "declare";
-      break;
-    case DbgVariableRecord::LocationType::Assign:
-      *OS << "assign";
-      break;
-    case DbgVariableRecord::LocationType::End:
-      *OS << "end";
-      break;
-    case DbgVariableRecord::LocationType::Any:
-      *OS << "any";
-      break;
-    };
+  void addPlugin(const VerifierPlugin *P) {
+    updatePlugins([&](List &Plugins) { Plugins.push_back(P); });
   }
 
-  void Write(const Metadata *MD) {
-    if (!MD)
-      return;
-    MD->print(*OS, MST, &M);
-    *OS << '\n';
+  void removePlugin(const VerifierPlugin *P) {
+    updatePlugins([&](List &Plugins) {
+      Plugins.erase(std::remove(Plugins.begin(), Plugins.end(), P),
+                    Plugins.end());
+    });
   }
 
-  template <class T> void Write(const MDTupleTypedArrayWrapper<T> &MD) {
-    Write(MD.get());
+  void addReader(PluginRegistryReader *R) {
+    if (llvm_is_multithreaded()) {
+      sys::ScopedLock Lock(Mutex);
+      Readers.push_back(R);
+    }
   }
 
-  void Write(const NamedMDNode *NMD) {
-    if (!NMD)
-      return;
-    NMD->print(*OS, MST);
-    *OS << '\n';
+  void removeReader(PluginRegistryReader *R) {
+    if (llvm_is_multithreaded()) {
+      sys::ScopedLock Lock(Mutex);
+      Readers.erase(std::remove(Readers.begin(), Readers.end(), R),
+                    Readers.end());
+    }
   }
+};
 
-  void Write(Type *T) {
-    if (!T)
-      return;
-    *OS << ' ' << *T;
-  }
+class PluginRegistryReader {
+  friend PluginRegistry;
+  friend PluginRegistryLock;
 
-  void Write(const Comdat *C) {
-    if (!C)
-      return;
-    *OS << *C;
-  }
+  std::atomic<PluginRegistry::List *> HazardPtr;
 
-  void Write(const APInt *AI) {
-    if (!AI)
-      return;
-    *OS << *AI << '\n';
+public:
+  PluginRegistryReader() {
+    HazardPtr.store(nullptr, std::memory_order_relaxed);
+    PluginRegistry::get().addReader(this);
   }
 
-  void Write(const unsigned i) { *OS << i << '\n'; }
+  ~PluginRegistryReader() { PluginRegistry::get().removeReader(this); }
+};
 
-  // NOLINTNEXTLINE(readability-identifier-naming)
-  void Write(const Attribute *A) {
-    if (!A)
-      return;
-    *OS << A->getAsString() << '\n';
-  }
+// Thread-safe update of the plugins list. Take the lock, copy & update the
+// list, then wait for all readers to let go of the old version of the list
+// before releasing the lock.
+template <typename FnT> void PluginRegistry::updatePlugins(FnT &&F) {
+  if (llvm_is_multithreaded()) {
+    sys::ScopedLock Lock(Mutex);
+
+    List *OldList = Plugins.load(std::memory_order_relaxed);
+    List *NewList = (OldList == &PluginsStorage[0]) ? &PluginsStorage[1]
+                                                    : &PluginsStorage[0];
+
+    // We're about to write to NewList. Spin wait to ensure no reader is
+    // accessing it.
+    for (auto *R : Readers) {
+      while (R->HazardPtr.load(std::memory_order_seq_cst) == NewList) {
+        // Let's yield to avoid a pathological busy wait. This really should
+        // only happen in the corner case where multiple users of LLVM exist
+        // in the same process and are initialized or torn down concurrently,
+        // so don't sweat the details.
+        std::this_thread::yield();
+      }
+    }
 
-  // NOLINTNEXTLINE(readability-identifier-naming)
-  void Write(const AttributeSet *AS) {
-    if (!AS)
-      return;
-    *OS << AS->getAsString() << '\n';
+    *NewList = *OldList;
+    F(*NewList);
+
+    Plugins.store(NewList, std::memory_order_seq_cst);
+  } else {
+    // Avoid unnecessary copies when compiling without multi-threading
+    // support.
+    F(*Plugins.load(std::memory_order_relaxed));
   }
+}
 
-  // NOLINTNEXTLINE(readability-identifier-naming)
-  void Write(const AttributeList *AL) {
-    if (!AL)
-      return;
-    AL->print(*OS);
+class PluginRegistryLock {
+  PluginRegistryLock(PluginRegistryLock &) = delete;
+  PluginRegistryLock(PluginRegistryLock &&) = delete;
+  PluginRegistryLock &operator=(PluginRegistryLock &) = delete;
+  PluginRegistryLock &operator=(PluginRegistryLock &&) = delete;
+
+  PluginRegistryReader &Reader;
+
+public:
+  explicit PluginRegistryLock(PluginRegistryReader &Reader) : Reader(Reader) {
+    assert(!Reader.HazardPtr &&
+           "cannot have multiple PluginRegistryLocks through the same reader");
+
+    auto &Registry = PluginRegistry::get();
+
+    // The memory order of the initial load is irrelevant since we re-check the
+    // pointer using a sequentially consistent load later.
+    PluginRegistry::List *L = Registry.Plugins.load(std::memory_order_relaxed);
+
+    if (llvm_is_multithreaded()) {
+      for (;;) {
+        Reader.HazardPtr.store(L, std::memory_order_seq_cst);
+
+        PluginRegistry::List *Check =
+            Registry.Plugins.load(std::memory_order_seq_cst);
+        if (Check == L)
+          break;
+
+        L = Check;
+      }
+    } else {
+      Reader.HazardPtr.store(L, std::memory_order_relaxed);
+    }
   }
 
-  void Write(Printable P) { *OS << P << '\n'; }
+  ~PluginRegistryLock() {
+    assert(Reader.HazardPtr);
 
-  template <typename T> void Write(ArrayRef<T> Vs) {
-    for (const T &V : Vs)
-      Write(V);
+    if (llvm_is_multithreaded()) {
+      Reader.HazardPtr.store(nullptr, std::memory_order_seq_cst);
+    } else {
+      Reader.HazardPtr.store(nullptr, std::memory_order_relaxed);
+    }
   }
 
-  template <typename T1, typename... Ts>
-  void WriteTs(const T1 &V1, const Ts &... Vs) {
-    Write(V1);
-    WriteTs(Vs...);
+  ArrayRef<const VerifierPlugin *> get() const {
+    return *Reader.HazardPtr.load(std::memory_order_relaxed);
   }
+};
 
-  template <typename... Ts> void WriteTs() {}
+} // anonymous namespace
 
-public:
-  /// A check failed, so printout out the condition and the message.
-  ///
-  /// This provides a nice place to put a breakpoint if you want to see why
-  /// something is not correct.
-  void CheckFailed(const Twine &Message) {
-    if (OS)
-      *OS << Message << '\n';
-    Broken = true;
+VerifierPlugin::VerifierPlugin() { PluginRegistry::get().addPlugin(this); }
+
+VerifierPlugin::~VerifierPlugin() { PluginRegistry::get().removePlugin(this); }
+
+void VerifierPlugin::verifyIntrinsicCall(CallBase &Call,
+                                         VerifierSupport &VS) const {}
+
+VerifierSupport::VerifierSupport(raw_ostream *OS, const Module &M)
+    : OS(OS), M(M), MST(&M), TT(M.getTargetTriple()), DL(M.getDataLayout()),
+      Context(M.getContext()) {}
+
+void VerifierSupport::Write(const Module *M) {
+  *OS << "; ModuleID = '" << M->getModuleIdentifier() << "'\n";
+}
+
+void VerifierSupport::Write(const Value *V) {
+  if (V)
+    Write(*V);
+}
+
+void VerifierSupport::Write(const Value &V) {
+  if (isa<Instruction>(V)) {
+    V.print(*OS, MST);
+    *OS << '\n';
+  } else {
+    V.printAsOperand(*OS, true, MST);
+    *OS << '\n';
   }
+}
 
-  /// A check failed (with values to print).
-  ///
-  /// This calls the Message-only version so that the above is easier to set a
-  /// breakpoint on.
-  template <typename T1, typename... Ts>
-  void CheckFailed(const Twine &Message, const T1 &V1, const Ts &... Vs) {
-    CheckFailed(Message);
-    if (OS)
-      WriteTs(V1, Vs...);
-  }
-
-  /// A debug info check failed.
-  void DebugInfoCheckFailed(const Twine &Message) {
-    if (OS)
-      *OS << Message << '\n';
-    Broken |= TreatBrokenDebugInfoAsError;
-    BrokenDebugInfo = true;
-  }
-
-  /// A debug info check failed (with values to print).
-  template <typename T1, typename... Ts>
-  void DebugInfoCheckFailed(const Twine &Message, const T1 &V1,
-                            const Ts &... Vs) {
-    DebugInfoCheckFailed(Message);
-    if (OS)
-      WriteTs(V1, Vs...);
+void VerifierSupport::Write(const DbgRecord *DR) {
+  if (DR) {
+    DR->print(*OS, MST, false);
+    *OS << '\n';
   }
-};
+}
+
+void VerifierSupport::Write(DbgVariableRecord::LocationType Type) {
+  switch (Type) {
+  case DbgVariableRecord::LocationType::Value:
+    *OS << "value";
+    break;
+  case DbgVariableRecord::LocationType::Declare:
+    *OS << "declare";
+    break;
+  case DbgVariableRecord::LocationType::Assign:
+    *OS << "assign";
+    break;
+  case DbgVariableRecord::LocationType::End:
+    *OS << "end";
+    break;
+  case DbgVariableRecord::LocationType::Any:
+    *OS << "any";
+    break;
+  };
+}
 
-} // namespace llvm
+void VerifierSupport::Write(const Metadata *MD) {
+  if (!MD)
+    return;
+  MD->print(*OS, MST, &M);
+  *OS << '\n';
+}
+
+void VerifierSupport::Write(const NamedMDNode *NMD) {
+  if (!NMD)
+    return;
+  NMD->print(*OS, MST);
+  *OS << '\n';
+}
+
+void VerifierSupport::Write(Type *T) {
+  if (!T)
+    return;
+  *OS << ' ' << *T;
+}
+
+void VerifierSupport::Write(const Comdat *C) {
+  if (!C)
+    return;
+  *OS << *C;
+}
+
+void VerifierSupport::Write(const APInt *AI) {
+  if (!AI)
+    return;
+  *OS << *AI << '\n';
+}
+
+// NOLINTNEXTLINE(readability-identifier-naming)
+void VerifierSupport::Write(const Attribute *A) {
+  if (!A)
+    return;
+  *OS << A->getAsString() << '\n';
+}
+
+// NOLINTNEXTLINE(readability-identifier-naming)
+void VerifierSupport::Write(const AttributeSet *AS) {
+  if (!AS)
+    return;
+  *OS << AS->getAsString() << '\n';
+}
+
+// NOLINTNEXTLINE(readability-identifier-naming)
+void VerifierSupport::Write(const AttributeList *AL) {
+  if (!AL)
+    return;
+  AL->print(*OS);
+}
+
+void VerifierSupport::CheckFailed(const Twine &Message) {
+  if (OS)
+    *OS << Message << '\n';
+  Broken = true;
+}
+
+/// A debug info check failed.
+void VerifierSupport::DebugInfoCheckFailed(const Twine &Message) {
+  if (OS)
+    *OS << Message << '\n';
+  Broken |= TreatBrokenDebugInfoAsError;
+  BrokenDebugInfo = true;
+}
 
 namespace {
 
 class Verifier : public InstVisitor<Verifier>, VerifierSupport {
   friend class InstVisitor<Verifier>;
+
+  PluginRegistryReader PluginsReader;
+
   DominatorTree DT;
 
   /// When verifying a basic block, keep track of all of the
@@ -5660,8 +5775,12 @@ void Verifier::visitIntrinsicCall(Intrinsic::ID ID, CallBase &Call) {
   }
 
   switch (ID) {
-  default:
+  default: {
+    PluginRegistryLock Lock(PluginsReader);
+    for (const VerifierPlugin *P : Lock.get())
+      P->verifyIntrinsicCall(Call, *this);
     break;
+  }
   case Intrinsic::assume: {
     for (auto &Elem : Call.bundle_op_infos()) {
       unsigned ArgCount = Elem.End - Elem.Begin;
@@ -6549,37 +6668,12 @@ void Verifier::visitIntrinsicCall(Intrinsic::ID ID, CallBase &Call) {
     break;
   }
   case Intrinsic::preserve_array_access_index:
-  case Intrinsic::preserve_struct_access_index:
-  case Intrinsic::aarch64_ldaxr:
-  case Intrinsic::aarch64_ldxr:
-  case Intrinsic::arm_ldaex:
-  case Intrinsic::arm_ldrex: {
+  case Intrinsic::preserve_struct_access_index: {
     Type *ElemTy = Call.getParamElementType(0);
     Check(ElemTy, "Intrinsic requires elementtype attribute on first argument.",
           &Call);
     break;
   }
-  case Intrinsic::aarch64_stlxr:
-  case Intrinsic::aarch64_stxr:
-  case Intrinsic::arm_stlex:
-  case Intrinsic::arm_strex: {
-    Type *ElemTy = Call.getAttributes().getParamElementType(1);
-    Check(ElemTy,
-          "Intrinsic requires elementtype attribute on second argument.",
-          &Call);
-    break;
-  }
-  case Intrinsic::aarch64_prefetch: {
-    Check(cast<ConstantInt>(Call.getArgOperand(1))->getZExtValue() < 2,
-          "write argument to llvm.aarch64.prefetch must be 0 or 1", Call);
-    Check(cast<ConstantInt>(Call.getArgOperand(2))->getZExtValue() < 4,
-          "target argument to llvm.aarch64.prefetch must be 0-3", Call);
-    Check(cast<ConstantInt>(Call.getArgOperand(3))->getZExtValue() < 2,
-          "stream argument to llvm.aarch64.prefetch must be 0 or 1", Call);
-    Check(cast<ConstantInt>(Call.getArgOperand(4))->getZExtValue() < 2,
-          "isdata argument to llvm.aarch64.prefetch must be 0 or 1", Call);
-    break;
-  }
   case Intrinsic::callbr_landingpad: {
     const auto *CBR = dyn_cast<CallBrInst>(Call.getOperand(0));
     Check(CBR, "intrinstic requires callbr operand", &Call);
@@ -6606,232 +6700,6 @@ void Verifier::visitIntrinsicCall(Intrinsic::ID ID, CallBase &Call) {
           &Call);
     break;
   }
-  case Intrinsic::amdgcn_cs_chain: {
-    auto CallerCC = Call.getCaller()->getCallingConv();
-    switch (CallerCC) {
-    case CallingConv::AMDGPU_CS:
-    case CallingConv::AMDGPU_CS_Chain:
-    case CallingConv::AMDGPU_CS_ChainPreserve:
-      break;
-    default:
-      CheckFailed("Intrinsic can only be used from functions with the "
-                  "amdgpu_cs, amdgpu_cs_chain or amdgpu_cs_chain_preserve "
-                  "calling conventions",
-                  &Call);
-      break;
-    }
-
-    Check(Call.paramHasAttr(2, Attribute::InReg),
-          "SGPR arguments must have the `inreg` attribute", &Call);
-    Check(!Call.paramHasAttr(3, Attribute::InReg),
-          "VGPR arguments must not have the `inreg` attribute", &Call);
-
-    auto *Next = Call.getNextNode();
-    bool IsAMDUnreachable = Next && isa<IntrinsicInst>(Next) &&
-                            cast<IntrinsicInst>(Next)->getIntrinsicID() ==
-                                Intrinsic::amdgcn_unreachable;
-    Check(Next && (isa<UnreachableInst>(Next) || IsAMDUnreachable),
-          "llvm.amdgcn.cs.chain must be followed by unreachable", &Call);
-    break;
-  }
-  case Intrinsic::amdgcn_init_exec_from_input: {
-    const Argument *Arg = dyn_cast<Argument>(Call.getOperand(0));
-    Check(Arg && Arg->hasInRegAttr(),
-          "only inreg arguments to the parent function are valid as inputs to "
-          "this intrinsic",
-          &Call);
-    break;
-  }
-  case Intrinsic::amdgcn_set_inactive_chain_arg: {
-    auto CallerCC = Call.getCaller()->getCallingConv();
-    switch (CallerCC) {
-    case CallingConv::AMDGPU_CS_Chain:
-    case CallingConv::AMDGPU_CS_ChainPreserve:
-      break;
-    default:
-      CheckFailed("Intrinsic can only be used from functions with the "
-                  "amdgpu_cs_chain or amdgpu_cs_chain_preserve "
-                  "calling conventions",
-                  &Call);
-      break;
-    }
-
-    unsigned InactiveIdx = 1;
-    Check(!Call.paramHasAttr(InactiveIdx, Attribute::InReg),
-          "Value for inactive lanes must not have the `inreg` attribute",
-          &Call);
-    Check(isa<Argument>(Call.getArgOperand(InactiveIdx)),
-          "Value for inactive lanes must be a function argument", &Call);
-    Check(!cast<Argument>(Call.getArgOperand(InactiveIdx))->hasInRegAttr(),
-          "Value for inactive lanes must be a VGPR function argument", &Call);
-    break;
-  }
-  case Intrinsic::amdgcn_call_whole_wave: {
-    auto F = dyn_cast<Function>(Call.getArgOperand(0));
-    Check(F, "Indirect whole wave calls are not allowed", &Call);
-
-    CallingConv::ID CC = F->getCallingConv();
-    Check(CC == CallingConv::AMDGPU_Gfx_WholeWave,
-          "Callee must have the amdgpu_gfx_whole_wave calling convention",
-          &Call);
-
-    Check(!F->isVarArg(), "Variadic whole wave calls are not allowed", &Call);
-
-    Check(Call.arg_size() == F->arg_size(),
-          "Call argument count must match callee argument count", &Call);
-
-    // The first argument of the call is the callee, and the first argument of
-    // the callee is the active mask. The rest of the arguments must match.
-    Check(F->arg_begin()->getType()->isIntegerTy(1),
-          "Callee must have i1 as its first argument", &Call);
-    for (auto [CallArg, FuncArg] :
-         drop_begin(zip_equal(Call.args(), F->args()))) {
-      Check(CallArg->getType() == FuncArg.getType(),
-            "Argument types must match", &Call);
-
-      // Check that inreg attributes match between call site and function
-      Check(Call.paramHasAttr(FuncArg.getArgNo(), Attribute::InReg) ==
-                FuncArg.hasInRegAttr(),
-            "Argument inreg attributes must match", &Call);
-    }
-    break;
-  }
-  case Intrinsic::amdgcn_s_prefetch_data: {
-    Check(
-        AMDGPU::isFlatGlobalAddrSpace(
-            Call.getArgOperand(0)->getType()->getPointerAddressSpace()),
-        "llvm.amdgcn.s.prefetch.data only supports global or constant memory");
-    break;
-  }
-  case Intrinsic::amdgcn_mfma_scale_f32_16x16x128_f8f6f4:
-  case Intrinsic::amdgcn_mfma_scale_f32_32x32x64_f8f6f4: {
-    Value *Src0 = Call.getArgOperand(0);
-    Value *Src1 = Call.getArgOperand(1);
-
-    uint64_t CBSZ = cast<ConstantInt>(Call.getArgOperand(3))->getZExtValue();
-    uint64_t BLGP = cast<ConstantInt>(Call.getArgOperand(4))->getZExtValue();
-    Check(CBSZ <= 4, "invalid value for cbsz format", Call,
-          Call.getArgOperand(3));
-    Check(BLGP <= 4, "invalid value for blgp format", Call,
-          Call.getArgOperand(4));
-
-    // AMDGPU::MFMAScaleFormats values
-    auto getFormatNumRegs = [](unsigned FormatVal) {
-      switch (FormatVal) {
-      case 0:
-      case 1:
-        return 8u;
-      case 2:
-      case 3:
-        return 6u;
-      case 4:
-        return 4u;
-      default:
-        llvm_unreachable("invalid format value");
-      }
-    };
-
-    auto isValidSrcASrcBVector = [](FixedVectorType *Ty) {
-      if (!Ty || !Ty->getElementType()->isIntegerTy(32))
-        return false;
-      unsigned NumElts = Ty->getNumElements();
-      return NumElts == 4 || NumElts == 6 || NumElts == 8;
-    };
-
-    auto *Src0Ty = dyn_cast<FixedVectorType>(Src0->getType());
-    auto *Src1Ty = dyn_cast<FixedVectorType>(Src1->getType());
-    Check(isValidSrcASrcBVector(Src0Ty),
-          "operand 0 must be 4, 6 or 8 element i32 vector", &Call, Src0);
-    Check(isValidSrcASrcBVector(Src1Ty),
-          "operand 1 must be 4, 6 or 8 element i32 vector", &Call, Src1);
-
-    // Permit excess registers for the format.
-    Check(Src0Ty->getNumElements() >= getFormatNumRegs(CBSZ),
-          "invalid vector type for format", &Call, Src0, Call.getArgOperand(3));
-    Check(Src1Ty->getNumElements() >= getFormatNumRegs(BLGP),
-          "invalid vector type for format", &Call, Src1, Call.getArgOperand(5));
-    break;
-  }
-  case Intrinsic::amdgcn_wmma_f32_16x16x128_f8f6f4:
-  case Intrinsic::amdgcn_wmma_scale_f32_16x16x128_f8f6f4:
-  case Intrinsic::amdgcn_wmma_scale16_f32_16x16x128_f8f6f4: {
-    Value *Src0 = Call.getArgOperand(1);
-    Value *Src1 = Call.getArgOperand(3);
-
-    unsigned FmtA = cast<ConstantInt>(Call.getArgOperand(0))->getZExtValue();
-    unsigned FmtB = cast<ConstantInt>(Call.getArgOperand(2))->getZExtValue();
-    Check(FmtA <= 4, "invalid value for matrix format", Call,
-          Call.getArgOperand(0));
-    Check(FmtB <= 4, "invalid value for matrix format", Call,
-          Call.getArgOperand(2));
-
-    // AMDGPU::MatrixFMT values
-    auto getFormatNumRegs = [](unsigned FormatVal) {
-      switch (FormatVal) {
-      case 0:
-      case 1:
-        return 16u;
-      case 2:
-      case 3:
-        return 12u;
-      case 4:
-        return 8u;
-      default:
-        llvm_unreachable("invalid format value");
-      }
-    };
-
-    auto isValidSrcASrcBVector = [](FixedVectorType *Ty) {
-      if (!Ty || !Ty->getElementType()->isIntegerTy(32))
-        return false;
-      unsigned NumElts = Ty->getNumElements();
-      return NumElts == 16 || NumElts == 12 || NumElts == 8;
-    };
-
-    auto *Src0Ty = dyn_cast<FixedVectorType>(Src0->getType());
-    auto *Src1Ty = dyn_cast<FixedVectorType>(Src1->getType());
-    Check(isValidSrcASrcBVector(Src0Ty),
-          "operand 1 must be 8, 12 or 16 element i32 vector", &Call, Src0);
-    Check(isValidSrcASrcBVector(Src1Ty),
-          "operand 3 must be 8, 12 or 16 element i32 vector", &Call, Src1);
-
-    // Permit excess registers for the format.
-    Check(Src0Ty->getNumElements() >= getFormatNumRegs(FmtA),
-          "invalid vector type for format", &Call, Src0, Call.getArgOperand(0));
-    Check(Src1Ty->getNumElements() >= getFormatNumRegs(FmtB),
-          "invalid vector type for format", &Call, Src1, Call.getArgOperand(2));
-    break;
-  }
-  case Intrinsic::amdgcn_cooperative_atomic_load_32x4B:
-  case Intrinsic::amdgcn_cooperative_atomic_load_16x8B:
-  case Intrinsic::amdgcn_cooperative_atomic_load_8x16B:
-  case Intrinsic::amdgcn_cooperative_atomic_store_32x4B:
-  case Intrinsic::amdgcn_cooperative_atomic_store_16x8B:
-  case Intrinsic::amdgcn_cooperative_atomic_store_8x16B: {
-    // Check we only use this intrinsic on the FLAT or GLOBAL address spaces.
-    Value *PtrArg = Call.getArgOperand(0);
-    const unsigned AS = PtrArg->getType()->getPointerAddressSpace();
-    Check(AS == AMDGPUAS::FLAT_ADDRESS || AS == AMDGPUAS::GLOBAL_ADDRESS,
-          "cooperative atomic intrinsics require a generic or global pointer",
-          &Call, PtrArg);
-
-    // Last argument must be a MD string
-    auto *Op = cast<MetadataAsValue>(Call.getArgOperand(Call.arg_size() - 1));
-    MDNode *MD = cast<MDNode>(Op->getMetadata());
-    Check((MD->getNumOperands() == 1) && isa<MDString>(MD->getOperand(0)),
-          "cooperative atomic intrinsics require that the last argument is a "
-          "metadata string",
-          &Call, Op);
-    break;
-  }
-  case Intrinsic::nvvm_setmaxnreg_inc_sync_aligned_u32:
-  case Intrinsic::nvvm_setmaxnreg_dec_sync_aligned_u32: {
-    Value *V = Call.getArgOperand(0);
-    unsigned RegCount = cast<ConstantInt>(V)->getZExtValue();
-    Check(RegCount % 8 == 0,
-          "reg_count argument to nvvm.setmaxnreg must be in multiples of 8");
-    break;
-  }
   case Intrinsic::experimental_convergence_entry:
   case Intrinsic::experimental_convergence_anchor:
     break;
diff --git a/llvm/lib/Target/AArch64/AArch64.h b/llvm/lib/Target/AArch64/AArch64.h
index 8d0ff41fc8c08..176159dd40766 100644
--- a/llvm/lib/Target/AArch64/AArch64.h
+++ b/llvm/lib/Target/AArch64/AArch64.h
@@ -115,6 +115,8 @@ void initializeSMEPeepholeOptPass(PassRegistry &);
 void initializeMachineSMEABIPass(PassRegistry &);
 void initializeSVEIntrinsicOptsPass(PassRegistry &);
 void initializeAArch64Arm64ECCallLoweringPass(PassRegistry &);
+
+void initializeAArch64Verifier();
 } // end namespace llvm
 
 #endif
diff --git a/llvm/lib/Target/AArch64/AArch64TargetMachine.cpp b/llvm/lib/Target/AArch64/AArch64TargetMachine.cpp
index dde1d88403bfe..122a4ebb39d84 100644
--- a/llvm/lib/Target/AArch64/AArch64TargetMachine.cpp
+++ b/llvm/lib/Target/AArch64/AArch64TargetMachine.cpp
@@ -239,6 +239,7 @@ LLVMInitializeAArch64Target() {
   RegisterTargetMachine<AArch64leTargetMachine> V(getTheAArch64_32Target());
   auto &PR = *PassRegistry::getPassRegistry();
   initializeGlobalISel(PR);
+  initializeAArch64Verifier();
   initializeAArch64A53Fix835769Pass(PR);
   initializeAArch64A57FPLoadBalancingPass(PR);
   initializeAArch64AdvSIMDScalarPass(PR);
diff --git a/llvm/lib/Target/AArch64/AArch64Verifier.cpp b/llvm/lib/Target/AArch64/AArch64Verifier.cpp
new file mode 100644
index 0000000000000..61e3618d11cc7
--- /dev/null
+++ b/llvm/lib/Target/AArch64/AArch64Verifier.cpp
@@ -0,0 +1,70 @@
+//===- AArch64Verifier.h --------------------------------------------------===//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+//
+/// IR verifier plugin for AArch64 intrinsics.
+//
+//===----------------------------------------------------------------------===//
+
+#include "AArch64.h"
+#include "llvm/IR/Constants.h"
+#include "llvm/IR/Instructions.h"
+#include "llvm/IR/IntrinsicsAArch64.h"
+#include "llvm/IR/Verifier.h"
+
+using namespace llvm;
+
+namespace {
+
+#define Check(C, ...)                                                          \
+  do {                                                                         \
+    if (!(C)) {                                                                \
+      VS.CheckFailed(__VA_ARGS__);                                             \
+      return;                                                                  \
+    }                                                                          \
+  } while (false)
+
+class AArch64Verifier : public VerifierPlugin {
+public:
+  void verifyIntrinsicCall(CallBase &Call, VerifierSupport &VS) const override {
+    switch (Call.getIntrinsicID()) {
+    default:
+      break;
+    case Intrinsic::aarch64_ldaxr:
+    case Intrinsic::aarch64_ldxr: {
+      Type *ElemTy = Call.getParamElementType(0);
+      Check(ElemTy,
+            "Intrinsic requires elementtype attribute on first argument.",
+            &Call);
+      break;
+    }
+    case Intrinsic::aarch64_stlxr:
+    case Intrinsic::aarch64_stxr: {
+      Type *ElemTy = Call.getParamElementType(1);
+      Check(ElemTy,
+            "Intrinsic requires elementtype attribute on second argument.",
+            &Call);
+      break;
+    }
+    case Intrinsic::aarch64_prefetch: {
+      Check(cast<ConstantInt>(Call.getArgOperand(1))->getZExtValue() < 2,
+            "write argument to llvm.aarch64.prefetch must be 0 or 1", Call);
+      Check(cast<ConstantInt>(Call.getArgOperand(2))->getZExtValue() < 4,
+            "target argument to llvm.aarch64.prefetch must be 0-3", Call);
+      Check(cast<ConstantInt>(Call.getArgOperand(3))->getZExtValue() < 2,
+            "stream argument to llvm.aarch64.prefetch must be 0 or 1", Call);
+      Check(cast<ConstantInt>(Call.getArgOperand(4))->getZExtValue() < 2,
+            "isdata argument to llvm.aarch64.prefetch must be 0 or 1", Call);
+      break;
+    }
+    }
+  }
+};
+
+} // anonymous namespace
+
+void llvm::initializeAArch64Verifier() { static AArch64Verifier Verifier; }
diff --git a/llvm/lib/Target/AArch64/CMakeLists.txt b/llvm/lib/Target/AArch64/CMakeLists.txt
index a8185358d6dfc..e96cc88b841d2 100644
--- a/llvm/lib/Target/AArch64/CMakeLists.txt
+++ b/llvm/lib/Target/AArch64/CMakeLists.txt
@@ -87,6 +87,7 @@ add_llvm_target(AArch64CodeGen
   AArch64TargetMachine.cpp
   AArch64TargetObjectFile.cpp
   AArch64TargetTransformInfo.cpp
+  AArch64Verifier.cpp
   SMEABIPass.cpp
   SMEPeepholeOpt.cpp
   SVEIntrinsicOpts.cpp
diff --git a/llvm/lib/Target/AMDGPU/AMDGPU.h b/llvm/lib/Target/AMDGPU/AMDGPU.h
index 0f2c33585884f..8aa4231a959e8 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPU.h
+++ b/llvm/lib/Target/AMDGPU/AMDGPU.h
@@ -22,6 +22,8 @@ class AMDGPUTargetMachine;
 class GCNTargetMachine;
 class TargetMachine;
 
+void initializeAMDGPUVerifier();
+
 // GlobalISel passes
 void initializeAMDGPUPreLegalizerCombinerPass(PassRegistry &);
 FunctionPass *createAMDGPUPreLegalizeCombiner(bool IsOptNone);
diff --git a/llvm/lib/Target/AMDGPU/AMDGPUTargetMachine.cpp b/llvm/lib/Target/AMDGPU/AMDGPUTargetMachine.cpp
index 92a587b5771b6..3a4a7a90911ae 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPUTargetMachine.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPUTargetMachine.cpp
@@ -531,6 +531,8 @@ extern "C" LLVM_ABI LLVM_EXTERNAL_VISIBILITY void LLVMInitializeAMDGPUTarget() {
   RegisterTargetMachine<R600TargetMachine> X(getTheR600Target());
   RegisterTargetMachine<GCNTargetMachine> Y(getTheGCNTarget());
 
+  initializeAMDGPUVerifier();
+
   PassRegistry *PR = PassRegistry::getPassRegistry();
   initializeR600ClauseMergePassPass(*PR);
   initializeR600ControlFlowFinalizerPass(*PR);
diff --git a/llvm/lib/Target/AMDGPU/AMDGPUVerifier.cpp b/llvm/lib/Target/AMDGPU/AMDGPUVerifier.cpp
new file mode 100644
index 0000000000000..5d555abd30a70
--- /dev/null
+++ b/llvm/lib/Target/AMDGPU/AMDGPUVerifier.cpp
@@ -0,0 +1,266 @@
+//===- AMDGPUVerifier.h ---------------------------------------------------===//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+//
+/// IR verifier plugin for AMDGPU intrinsics.
+//
+//===----------------------------------------------------------------------===//
+
+#include "AMDGPU.h"
+#include "llvm/IR/Instructions.h"
+#include "llvm/IR/IntrinsicInst.h"
+#include "llvm/IR/IntrinsicsAMDGPU.h"
+#include "llvm/IR/Verifier.h"
+
+using namespace llvm;
+
+namespace {
+
+#define Check(C, ...)                                                          \
+  do {                                                                         \
+    if (!(C)) {                                                                \
+      VS.CheckFailed(__VA_ARGS__);                                             \
+      return;                                                                  \
+    }                                                                          \
+  } while (false)
+
+class AMDGPUVerifier : public VerifierPlugin {
+public:
+  void verifyIntrinsicCall(CallBase &Call, VerifierSupport &VS) const override {
+    switch (Call.getIntrinsicID()) {
+    default:
+      break;
+    case Intrinsic::amdgcn_cs_chain: {
+      auto CallerCC = Call.getCaller()->getCallingConv();
+      switch (CallerCC) {
+      case CallingConv::AMDGPU_CS:
+      case CallingConv::AMDGPU_CS_Chain:
+      case CallingConv::AMDGPU_CS_ChainPreserve:
+        break;
+      default:
+        VS.CheckFailed("Intrinsic can only be used from functions with the "
+                       "amdgpu_cs, amdgpu_cs_chain or amdgpu_cs_chain_preserve "
+                       "calling conventions",
+                       &Call);
+        break;
+      }
+
+      Check(Call.paramHasAttr(2, Attribute::InReg),
+            "SGPR arguments must have the `inreg` attribute", &Call);
+      Check(!Call.paramHasAttr(3, Attribute::InReg),
+            "VGPR arguments must not have the `inreg` attribute", &Call);
+
+      auto *Next = Call.getNextNode();
+      bool IsAMDUnreachable = Next && isa<IntrinsicInst>(Next) &&
+                              cast<IntrinsicInst>(Next)->getIntrinsicID() ==
+                                  Intrinsic::amdgcn_unreachable;
+      Check(Next && (isa<UnreachableInst>(Next) || IsAMDUnreachable),
+            "llvm.amdgcn.cs.chain must be followed by unreachable", &Call);
+      break;
+    }
+    case Intrinsic::amdgcn_init_exec_from_input: {
+      const Argument *Arg = dyn_cast<Argument>(Call.getOperand(0));
+      Check(
+          Arg && Arg->hasInRegAttr(),
+          "only inreg arguments to the parent function are valid as inputs to "
+          "this intrinsic",
+          &Call);
+      break;
+    }
+    case Intrinsic::amdgcn_set_inactive_chain_arg: {
+      auto CallerCC = Call.getCaller()->getCallingConv();
+      switch (CallerCC) {
+      case CallingConv::AMDGPU_CS_Chain:
+      case CallingConv::AMDGPU_CS_ChainPreserve:
+        break;
+      default:
+        VS.CheckFailed("Intrinsic can only be used from functions with the "
+                       "amdgpu_cs_chain or amdgpu_cs_chain_preserve "
+                       "calling conventions",
+                       &Call);
+        break;
+      }
+
+      unsigned InactiveIdx = 1;
+      Check(!Call.paramHasAttr(InactiveIdx, Attribute::InReg),
+            "Value for inactive lanes must not have the `inreg` attribute",
+            &Call);
+      Check(isa<Argument>(Call.getArgOperand(InactiveIdx)),
+            "Value for inactive lanes must be a function argument", &Call);
+      Check(!cast<Argument>(Call.getArgOperand(InactiveIdx))->hasInRegAttr(),
+            "Value for inactive lanes must be a VGPR function argument", &Call);
+      break;
+    }
+    case Intrinsic::amdgcn_call_whole_wave: {
+      auto F = dyn_cast<Function>(Call.getArgOperand(0));
+      Check(F, "Indirect whole wave calls are not allowed", &Call);
+
+      CallingConv::ID CC = F->getCallingConv();
+      Check(CC == CallingConv::AMDGPU_Gfx_WholeWave,
+            "Callee must have the amdgpu_gfx_whole_wave calling convention",
+            &Call);
+
+      Check(!F->isVarArg(), "Variadic whole wave calls are not allowed", &Call);
+
+      Check(Call.arg_size() == F->arg_size(),
+            "Call argument count must match callee argument count", &Call);
+
+      // The first argument of the call is the callee, and the first argument of
+      // the callee is the active mask. The rest of the arguments must match.
+      Check(F->arg_begin()->getType()->isIntegerTy(1),
+            "Callee must have i1 as its first argument", &Call);
+      for (auto [CallArg, FuncArg] :
+           drop_begin(zip_equal(Call.args(), F->args()))) {
+        Check(CallArg->getType() == FuncArg.getType(),
+              "Argument types must match", &Call);
+
+        // Check that inreg attributes match between call site and function
+        Check(Call.paramHasAttr(FuncArg.getArgNo(), Attribute::InReg) ==
+                  FuncArg.hasInRegAttr(),
+              "Argument inreg attributes must match", &Call);
+      }
+      break;
+    }
+    case Intrinsic::amdgcn_s_prefetch_data: {
+      Check(AMDGPU::isFlatGlobalAddrSpace(
+                Call.getArgOperand(0)->getType()->getPointerAddressSpace()),
+            "llvm.amdgcn.s.prefetch.data only supports global or constant "
+            "memory");
+      break;
+    }
+    case Intrinsic::amdgcn_mfma_scale_f32_16x16x128_f8f6f4:
+    case Intrinsic::amdgcn_mfma_scale_f32_32x32x64_f8f6f4: {
+      Value *Src0 = Call.getArgOperand(0);
+      Value *Src1 = Call.getArgOperand(1);
+
+      uint64_t CBSZ = cast<ConstantInt>(Call.getArgOperand(3))->getZExtValue();
+      uint64_t BLGP = cast<ConstantInt>(Call.getArgOperand(4))->getZExtValue();
+      Check(CBSZ <= 4, "invalid value for cbsz format", Call,
+            Call.getArgOperand(3));
+      Check(BLGP <= 4, "invalid value for blgp format", Call,
+            Call.getArgOperand(4));
+
+      // AMDGPU::MFMAScaleFormats values
+      auto getFormatNumRegs = [](unsigned FormatVal) {
+        switch (FormatVal) {
+        case 0:
+        case 1:
+          return 8u;
+        case 2:
+        case 3:
+          return 6u;
+        case 4:
+          return 4u;
+        default:
+          llvm_unreachable("invalid format value");
+        }
+      };
+
+      auto isValidSrcASrcBVector = [](FixedVectorType *Ty) {
+        if (!Ty || !Ty->getElementType()->isIntegerTy(32))
+          return false;
+        unsigned NumElts = Ty->getNumElements();
+        return NumElts == 4 || NumElts == 6 || NumElts == 8;
+      };
+
+      auto *Src0Ty = dyn_cast<FixedVectorType>(Src0->getType());
+      auto *Src1Ty = dyn_cast<FixedVectorType>(Src1->getType());
+      Check(isValidSrcASrcBVector(Src0Ty),
+            "operand 0 must be 4, 6 or 8 element i32 vector", &Call, Src0);
+      Check(isValidSrcASrcBVector(Src1Ty),
+            "operand 1 must be 4, 6 or 8 element i32 vector", &Call, Src1);
+
+      // Permit excess registers for the format.
+      Check(Src0Ty->getNumElements() >= getFormatNumRegs(CBSZ),
+            "invalid vector type for format", &Call, Src0,
+            Call.getArgOperand(3));
+      Check(Src1Ty->getNumElements() >= getFormatNumRegs(BLGP),
+            "invalid vector type for format", &Call, Src1,
+            Call.getArgOperand(5));
+      break;
+    }
+    case Intrinsic::amdgcn_wmma_f32_16x16x128_f8f6f4:
+    case Intrinsic::amdgcn_wmma_scale_f32_16x16x128_f8f6f4:
+    case Intrinsic::amdgcn_wmma_scale16_f32_16x16x128_f8f6f4: {
+      Value *Src0 = Call.getArgOperand(1);
+      Value *Src1 = Call.getArgOperand(3);
+
+      unsigned FmtA = cast<ConstantInt>(Call.getArgOperand(0))->getZExtValue();
+      unsigned FmtB = cast<ConstantInt>(Call.getArgOperand(2))->getZExtValue();
+      Check(FmtA <= 4, "invalid value for matrix format", Call,
+            Call.getArgOperand(0));
+      Check(FmtB <= 4, "invalid value for matrix format", Call,
+            Call.getArgOperand(2));
+
+      // AMDGPU::MatrixFMT values
+      auto getFormatNumRegs = [](unsigned FormatVal) {
+        switch (FormatVal) {
+        case 0:
+        case 1:
+          return 16u;
+        case 2:
+        case 3:
+          return 12u;
+        case 4:
+          return 8u;
+        default:
+          llvm_unreachable("invalid format value");
+        }
+      };
+
+      auto isValidSrcASrcBVector = [](FixedVectorType *Ty) {
+        if (!Ty || !Ty->getElementType()->isIntegerTy(32))
+          return false;
+        unsigned NumElts = Ty->getNumElements();
+        return NumElts == 16 || NumElts == 12 || NumElts == 8;
+      };
+
+      auto *Src0Ty = dyn_cast<FixedVectorType>(Src0->getType());
+      auto *Src1Ty = dyn_cast<FixedVectorType>(Src1->getType());
+      Check(isValidSrcASrcBVector(Src0Ty),
+            "operand 1 must be 8, 12 or 16 element i32 vector", &Call, Src0);
+      Check(isValidSrcASrcBVector(Src1Ty),
+            "operand 3 must be 8, 12 or 16 element i32 vector", &Call, Src1);
+
+      // Permit excess registers for the format.
+      Check(Src0Ty->getNumElements() >= getFormatNumRegs(FmtA),
+            "invalid vector type for format", &Call, Src0,
+            Call.getArgOperand(0));
+      Check(Src1Ty->getNumElements() >= getFormatNumRegs(FmtB),
+            "invalid vector type for format", &Call, Src1,
+            Call.getArgOperand(2));
+      break;
+    }
+    case Intrinsic::amdgcn_cooperative_atomic_load_32x4B:
+    case Intrinsic::amdgcn_cooperative_atomic_load_16x8B:
+    case Intrinsic::amdgcn_cooperative_atomic_load_8x16B:
+    case Intrinsic::amdgcn_cooperative_atomic_store_32x4B:
+    case Intrinsic::amdgcn_cooperative_atomic_store_16x8B:
+    case Intrinsic::amdgcn_cooperative_atomic_store_8x16B: {
+      // Check we only use this intrinsic on the FLAT or GLOBAL address spaces.
+      Value *PtrArg = Call.getArgOperand(0);
+      const unsigned AS = PtrArg->getType()->getPointerAddressSpace();
+      Check(AS == AMDGPUAS::FLAT_ADDRESS || AS == AMDGPUAS::GLOBAL_ADDRESS,
+            "cooperative atomic intrinsics require a generic or global pointer",
+            &Call, PtrArg);
+
+      // Last argument must be a MD string
+      auto *Op = cast<MetadataAsValue>(Call.getArgOperand(Call.arg_size() - 1));
+      MDNode *MD = cast<MDNode>(Op->getMetadata());
+      Check((MD->getNumOperands() == 1) && isa<MDString>(MD->getOperand(0)),
+            "cooperative atomic intrinsics require that the last argument is a "
+            "metadata string",
+            &Call, Op);
+      break;
+    }
+    }
+  }
+};
+
+} // anonymous namespace
+
+void llvm::initializeAMDGPUVerifier() { static AMDGPUVerifier TheVerifier; }
diff --git a/llvm/lib/Target/AMDGPU/CMakeLists.txt b/llvm/lib/Target/AMDGPU/CMakeLists.txt
index aae56eef73edd..2b04cdf917412 100644
--- a/llvm/lib/Target/AMDGPU/CMakeLists.txt
+++ b/llvm/lib/Target/AMDGPU/CMakeLists.txt
@@ -117,6 +117,7 @@ add_llvm_target(AMDGPUCodeGen
   AMDGPUTargetMachine.cpp
   AMDGPUTargetObjectFile.cpp
   AMDGPUTargetTransformInfo.cpp
+  AMDGPUVerifier.cpp
   AMDGPUWaitSGPRHazards.cpp
   AMDGPUUnifyDivergentExitNodes.cpp
   R600MachineCFGStructurizer.cpp
diff --git a/llvm/lib/Target/ARM/ARM.h b/llvm/lib/Target/ARM/ARM.h
index 3847f4e966afe..e34ab59683faa 100644
--- a/llvm/lib/Target/ARM/ARM.h
+++ b/llvm/lib/Target/ARM/ARM.h
@@ -81,6 +81,8 @@ void initializeMVEVPTBlockPass(PassRegistry &);
 void initializeThumb2ITBlockPass(PassRegistry &);
 void initializeThumb2SizeReducePass(PassRegistry &);
 
+void initializeARMVerifier();
+
 } // end namespace llvm
 
 #endif // LLVM_LIB_TARGET_ARM_ARM_H
diff --git a/llvm/lib/Target/ARM/ARMTargetMachine.cpp b/llvm/lib/Target/ARM/ARMTargetMachine.cpp
index 346776e0c4b25..3e6b9688b933d 100644
--- a/llvm/lib/Target/ARM/ARMTargetMachine.cpp
+++ b/llvm/lib/Target/ARM/ARMTargetMachine.cpp
@@ -90,6 +90,8 @@ extern "C" LLVM_ABI LLVM_EXTERNAL_VISIBILITY void LLVMInitializeARMTarget() {
   RegisterTargetMachine<ARMBETargetMachine> Y(getTheARMBETarget());
   RegisterTargetMachine<ARMBETargetMachine> B(getTheThumbBETarget());
 
+  initializeARMVerifier();
+
   PassRegistry &Registry = *PassRegistry::getPassRegistry();
   initializeGlobalISel(Registry);
   initializeARMAsmPrinterPass(Registry);
diff --git a/llvm/lib/Target/ARM/ARMVerifier.cpp b/llvm/lib/Target/ARM/ARMVerifier.cpp
new file mode 100644
index 0000000000000..a4e116fb64aef
--- /dev/null
+++ b/llvm/lib/Target/ARM/ARMVerifier.cpp
@@ -0,0 +1,58 @@
+//===- ARMVerifier.h ------------------------------------------------------===//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+//
+/// IR verifier plugin for ARM intrinsics.
+//
+//===----------------------------------------------------------------------===//
+
+#include "ARM.h"
+#include "llvm/IR/Instructions.h"
+#include "llvm/IR/IntrinsicsARM.h"
+#include "llvm/IR/Verifier.h"
+
+using namespace llvm;
+
+namespace {
+
+#define Check(C, ...)                                                          \
+  do {                                                                         \
+    if (!(C)) {                                                                \
+      VS.CheckFailed(__VA_ARGS__);                                             \
+      return;                                                                  \
+    }                                                                          \
+  } while (false)
+
+class ARMVerifier : public VerifierPlugin {
+public:
+  void verifyIntrinsicCall(CallBase &Call, VerifierSupport &VS) const override {
+    switch (Call.getIntrinsicID()) {
+    default:
+      break;
+    case Intrinsic::arm_ldaex:
+    case Intrinsic::arm_ldrex: {
+      Type *ElemTy = Call.getParamElementType(0);
+      Check(ElemTy,
+            "Intrinsic requires elementtype attribute on first argument.",
+            &Call);
+      break;
+    }
+    case Intrinsic::arm_stlex:
+    case Intrinsic::arm_strex: {
+      Type *ElemTy = Call.getParamElementType(1);
+      Check(ElemTy,
+            "Intrinsic requires elementtype attribute on second argument.",
+            &Call);
+      break;
+    }
+    }
+  }
+};
+
+} // anonymous namespace
+
+void llvm::initializeARMVerifier() { static ARMVerifier Verifier; }
diff --git a/llvm/lib/Target/ARM/CMakeLists.txt b/llvm/lib/Target/ARM/CMakeLists.txt
index fa778cad4af8e..60d28421486c7 100644
--- a/llvm/lib/Target/ARM/CMakeLists.txt
+++ b/llvm/lib/Target/ARM/CMakeLists.txt
@@ -57,6 +57,7 @@ add_llvm_target(ARMCodeGen
   ARMTargetMachine.cpp
   ARMTargetObjectFile.cpp
   ARMTargetTransformInfo.cpp
+  ARMVerifier.cpp
   MLxExpansionPass.cpp
   MVEGatherScatterLowering.cpp
   MVELaneInterleavingPass.cpp
diff --git a/llvm/lib/Target/NVPTX/CMakeLists.txt b/llvm/lib/Target/NVPTX/CMakeLists.txt
index 693f0d0b35edc..ce8c3fbbd0e7a 100644
--- a/llvm/lib/Target/NVPTX/CMakeLists.txt
+++ b/llvm/lib/Target/NVPTX/CMakeLists.txt
@@ -42,6 +42,7 @@ set(NVPTXCodeGen_sources
   NVPTXUtilities.cpp
   NVVMIntrRange.cpp
   NVVMReflect.cpp
+  NVVMVerifier.cpp
   )
 
 add_llvm_target(NVPTXCodeGen
diff --git a/llvm/lib/Target/NVPTX/NVPTX.h b/llvm/lib/Target/NVPTX/NVPTX.h
index 77a0e03d4075a..cb8f9e15154be 100644
--- a/llvm/lib/Target/NVPTX/NVPTX.h
+++ b/llvm/lib/Target/NVPTX/NVPTX.h
@@ -56,6 +56,8 @@ MachineFunctionPass *createNVPTXPeephole();
 MachineFunctionPass *createNVPTXProxyRegErasurePass();
 MachineFunctionPass *createNVPTXForwardParamsPass();
 
+void initializeNVVMVerifier();
+
 void initializeNVVMReflectLegacyPassPass(PassRegistry &);
 void initializeGenericToNVVMLegacyPassPass(PassRegistry &);
 void initializeNVPTXAllocaHoistingPass(PassRegistry &);
diff --git a/llvm/lib/Target/NVPTX/NVPTXTargetMachine.cpp b/llvm/lib/Target/NVPTX/NVPTXTargetMachine.cpp
index a6837a482608c..789620d3e3b73 100644
--- a/llvm/lib/Target/NVPTX/NVPTXTargetMachine.cpp
+++ b/llvm/lib/Target/NVPTX/NVPTXTargetMachine.cpp
@@ -93,6 +93,8 @@ extern "C" LLVM_ABI LLVM_EXTERNAL_VISIBILITY void LLVMInitializeNVPTXTarget() {
   RegisterTargetMachine<NVPTXTargetMachine32> X(getTheNVPTXTarget32());
   RegisterTargetMachine<NVPTXTargetMachine64> Y(getTheNVPTXTarget64());
 
+  initializeNVVMVerifier();
+
   PassRegistry &PR = *PassRegistry::getPassRegistry();
   // FIXME: This pass is really intended to be invoked during IR optimization,
   // but it's very NVPTX-specific.
diff --git a/llvm/lib/Target/NVPTX/NVVMVerifier.cpp b/llvm/lib/Target/NVPTX/NVVMVerifier.cpp
new file mode 100644
index 0000000000000..7f8a471b76bfa
--- /dev/null
+++ b/llvm/lib/Target/NVPTX/NVVMVerifier.cpp
@@ -0,0 +1,51 @@
+//===- NVVMVerifier.h -----------------------------------------------------===//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+//
+/// IR verifier plugin for NVVM intrinsics.
+//
+//===----------------------------------------------------------------------===//
+
+#include "NVPTX.h"
+#include "llvm/IR/Constants.h"
+#include "llvm/IR/Instructions.h"
+#include "llvm/IR/IntrinsicsNVPTX.h"
+#include "llvm/IR/Verifier.h"
+
+using namespace llvm;
+
+namespace {
+
+#define Check(C, ...)                                                          \
+  do {                                                                         \
+    if (!(C)) {                                                                \
+      VS.CheckFailed(__VA_ARGS__);                                             \
+      return;                                                                  \
+    }                                                                          \
+  } while (false)
+
+class NVVMVerifier : public VerifierPlugin {
+public:
+  void verifyIntrinsicCall(CallBase &Call, VerifierSupport &VS) const override {
+    switch (Call.getIntrinsicID()) {
+    default:
+      break;
+    case Intrinsic::nvvm_setmaxnreg_inc_sync_aligned_u32:
+    case Intrinsic::nvvm_setmaxnreg_dec_sync_aligned_u32: {
+      Value *V = Call.getArgOperand(0);
+      unsigned RegCount = cast<ConstantInt>(V)->getZExtValue();
+      Check(RegCount % 8 == 0,
+            "reg_count argument to nvvm.setmaxnreg must be in multiples of 8");
+      break;
+    }
+    }
+  }
+};
+
+} // anonymous namespace
+
+void llvm::initializeNVVMVerifier() { static NVVMVerifier Verifier; }
diff --git a/llvm/tools/llvm-as/CMakeLists.txt b/llvm/tools/llvm-as/CMakeLists.txt
index b21410fd23af7..7ba9f7ee9fa55 100644
--- a/llvm/tools/llvm-as/CMakeLists.txt
+++ b/llvm/tools/llvm-as/CMakeLists.txt
@@ -1,4 +1,6 @@
 set(LLVM_LINK_COMPONENTS
+  AllTargetsCodeGens
+  AllTargetsInfos
   AsmParser
   BitWriter
   Core
diff --git a/llvm/tools/llvm-as/llvm-as.cpp b/llvm/tools/llvm-as/llvm-as.cpp
index 21648674b51f1..5f1fcbc981156 100644
--- a/llvm/tools/llvm-as/llvm-as.cpp
+++ b/llvm/tools/llvm-as/llvm-as.cpp
@@ -25,6 +25,7 @@
 #include "llvm/Support/InitLLVM.h"
 #include "llvm/Support/SourceMgr.h"
 #include "llvm/Support/SystemUtils.h"
+#include "llvm/Support/TargetSelect.h"
 #include "llvm/Support/ToolOutputFile.h"
 #include <memory>
 #include <optional>
@@ -114,6 +115,7 @@ static void WriteOutputFile(const Module *M, const ModuleSummaryIndex *Index) {
 
 int main(int argc, char **argv) {
   InitLLVM X(argc, argv);
+  InitializeAllTargets(); // for verifier plugins
   cl::HideUnrelatedOptions(AsCat);
   cl::ParseCommandLineOptions(argc, argv, "llvm .ll -> .bc assembler\n");
   LLVMContext Context;