[mlir] make implementations of translation to LLVM IR interfaces private

[ftynse]

Summary

There is no need for the interface implementations to be exposed, opaque
registration functions are sufficient for all users, similarly to passes.

flang/lib/Optimizer/CMakeLists.txt

add_flang_library(FIROptimizer
 
   DEPENDS
   FIROpsIncGen
   FIROptTransformsPassIncGen
   ${dialect_libs}
 
   LINK_LIBS
   ${dialect_libs}
-  MLIRTargetLLVMIR
-  MLIRTargetLLVMIRModuleTranslation
+  MLIRLLVMToLLVMIRTranslation
+  MLIRTargetLLVMIRExport
 
   LINK_COMPONENTS
   AsmParser
   AsmPrinter
   Remarks
 )

mlir/examples/toy/Ch6/CMakeLists.txt

PRIVATE
     ${dialect_libs}
     ${conversion_libs}
     MLIRAnalysis
     MLIRCallInterfaces
     MLIRCastInterfaces
     MLIRExecutionEngine
     MLIRIR
     MLIRLLVMIR
+    MLIRLLVMToLLVMIRTranslation
     MLIRParser
     MLIRPass
     MLIRSideEffectInterfaces
     MLIRSupport
-    MLIRTargetLLVMIR
+    MLIRTargetLLVMIRExport
     MLIRTransforms
     )

mlir/examples/toy/Ch6/toyc.cpp

 #include "mlir/IR/AsmState.h"
 #include "mlir/IR/BuiltinOps.h"
 #include "mlir/IR/MLIRContext.h"
 #include "mlir/IR/Verifier.h"
 #include "mlir/InitAllDialects.h"
 #include "mlir/Parser.h"
 #include "mlir/Pass/Pass.h"
 #include "mlir/Pass/PassManager.h"
-#include "mlir/Target/LLVMIR.h"
+#include "mlir/Target/LLVMIR/Dialect/LLVMIR/LLVMToLLVMIRTranslation.h"
 #include "mlir/Target/LLVMIR/Export.h"
 #include "mlir/Transforms/Passes.h"
 
 #include "llvm/ADT/StringRef.h"
 #include "llvm/IR/Module.h"
 #include "llvm/Support/CommandLine.h"
 #include "llvm/Support/ErrorOr.h"
 #include "llvm/Support/MemoryBuffer.h"

mlir/examples/toy/Ch7/CMakeLists.txt

target_link_libraries(toyc-ch7
   PRIVATE
     ${dialect_libs}
     ${conversion_libs}
     MLIRAnalysis
     MLIRCallInterfaces
     MLIRCastInterfaces
     MLIRExecutionEngine
     MLIRIR
+    MLIRLLVMToLLVMIRTranslation
     MLIRParser
     MLIRPass
     MLIRSideEffectInterfaces
-    MLIRTargetLLVMIR
+    MLIRTargetLLVMIRExport
     MLIRTransforms
     )

mlir/examples/toy/Ch7/toyc.cpp

 #include "mlir/IR/AsmState.h"
 #include "mlir/IR/BuiltinOps.h"
 #include "mlir/IR/MLIRContext.h"
 #include "mlir/IR/Verifier.h"
 #include "mlir/InitAllDialects.h"
 #include "mlir/Parser.h"
 #include "mlir/Pass/Pass.h"
 #include "mlir/Pass/PassManager.h"
-#include "mlir/Target/LLVMIR.h"
+#include "mlir/Target/LLVMIR/Dialect/LLVMIR/LLVMToLLVMIRTranslation.h"
 #include "mlir/Target/LLVMIR/Export.h"
 #include "mlir/Transforms/Passes.h"
 
 #include "llvm/ADT/StringRef.h"
 #include "llvm/IR/Module.h"
 #include "llvm/Support/CommandLine.h"
 #include "llvm/Support/ErrorOr.h"
 #include "llvm/Support/MemoryBuffer.h"

mlir/include/mlir/Target/LLVMIR.h

 /// Convert the given LLVM module into MLIR's LLVM dialect.  The LLVM context is
 /// extracted from the registered LLVM IR dialect. In case of error, report it
 /// to the error handler registered with the MLIR context, if any (obtained from
 /// the MLIR module), and return `{}`.
 OwningModuleRef
 translateLLVMIRToModule(std::unique_ptr<llvm::Module> llvmModule,
                         MLIRContext *context);
 
-/// Register the LLVM dialect and the translation from it to the LLVM IR in the
-/// given registry;
-void registerLLVMDialectTranslation(DialectRegistry &registry);
-
-/// Register the LLVM dialect and the translation from it in the registry
-/// associated with the given context. This checks if the interface is already
-/// registered and avoids double registation.
-void registerLLVMDialectTranslation(MLIRContext &context);
-
 } // namespace mlir
 
 #endif // MLIR_TARGET_LLVMIR_H

mlir/include/mlir/Target/LLVMIR/Dialect/All.h

+//===- All.h - MLIR To LLVM IR Translation Registration ---------*- C++ -*-===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+//
+// This file defines a helper to register the translations of all suitable
+// dialects to LLVM IR.
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef MLIR_TARGET_LLVMIR_DIALECT_ALL_H
+#define MLIR_TARGET_LLVMIR_DIALECT_ALL_H
+
+#include "mlir/Target/LLVMIR/Dialect/LLVMAVX512/LLVMAVX512ToLLVMIRTranslation.h"
+#include "mlir/Target/LLVMIR/Dialect/LLVMArmNeon/LLVMArmNeonToLLVMIRTranslation.h"
+#include "mlir/Target/LLVMIR/Dialect/LLVMArmSVE/LLVMArmSVEToLLVMIRTranslation.h"
+#include "mlir/Target/LLVMIR/Dialect/LLVMIR/LLVMToLLVMIRTranslation.h"
+#include "mlir/Target/LLVMIR/Dialect/NVVM/NVVMToLLVMIRTranslation.h"
+#include "mlir/Target/LLVMIR/Dialect/OpenMP/OpenMPToLLVMIRTranslation.h"
+#include "mlir/Target/LLVMIR/Dialect/ROCDL/ROCDLToLLVMIRTranslation.h"
+
+namespace mlir {
+class DialectRegistry;
+
+/// Registers all dialects that can be translated to LLVM IR and the
+/// corresponding translation interfaces.
+static inline void registerAllToLLVMIRTranslations(DialectRegistry &registry) {
+  registerLLVMAVX512DialectTranslation(registry);
+  registerLLVMArmNeonDialectTranslation(registry);
+  registerLLVMArmSVEDialectTranslation(registry);
+  registerLLVMDialectTranslation(registry);
+  registerNVVMDialectTranslation(registry);
+  registerOpenMPDialectTranslation(registry);
+  registerROCDLDialectTranslation(registry);
+}
+} // namespace mlir
+
+#endif // MLIR_TARGET_LLVMIR_DIALECT_ALL_H

mlir/include/mlir/Target/LLVMIR/Dialect/LLVMAVX512/LLVMAVX512ToLLVMIRTranslation.h

 //===- LLVMAVX512ToLLVMIRTranslation.h - LLVMAVX512 to LLVM IR --*- C++ -*-===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 //
-// This file implements the dialect interface for translating the LLVMAVX512
-// dialect to LLVM IR.
+// This provides registration calls for LLVMAVX512 dialect to LLVM IR
+// translation.
 //
 //===----------------------------------------------------------------------===//
 
 #ifndef MLIR_TARGET_LLVMIR_DIALECT_LLVMAVX512_LLVMAVX512TOLLVMIRTRANSLATION_H
 #define MLIR_TARGET_LLVMIR_DIALECT_LLVMAVX512_LLVMAVX512TOLLVMIRTRANSLATION_H
 
-#include "mlir/Target/LLVMIR/LLVMTranslationInterface.h"
-
 namespace mlir {
 
-/// Implementation of the dialect interface that converts operations belonging
-/// to the LLVMAVX512 dialect to LLVM IR.
-class LLVMAVX512DialectLLVMIRTranslationInterface
-    : public LLVMTranslationDialectInterface {
-public:
-  using LLVMTranslationDialectInterface::LLVMTranslationDialectInterface;
-
-  /// Translates the given operation to LLVM IR using the provided IR builder
-  /// and saving the state in `moduleTranslation`.
-  LogicalResult
-  convertOperation(Operation *op, llvm::IRBuilderBase &builder,
-                   LLVM::ModuleTranslation &moduleTranslation) const final;
-};
+class DialectRegistry;
+class MLIRContext;
+
+/// Register the LLVMAVX512 dialect and the translation from it to the LLVM IR
+/// in the given registry;
+void registerLLVMAVX512DialectTranslation(DialectRegistry &registry);
+
+/// Register the LLVMAVX512 dialect and the translation from it in the registry
+/// associated with the given context.
+void registerLLVMAVX512DialectTranslation(MLIRContext &context);
 
 } // namespace mlir
 
 #endif // MLIR_TARGET_LLVMIR_DIALECT_LLVMAVX512_LLVMAVX512TOLLVMIRTRANSLATION_H

mlir/include/mlir/Target/LLVMIR/Dialect/LLVMArmNeon/LLVMArmNeonToLLVMIRTranslation.h

 //===- LLVMArmNeonToLLVMIRTranslation.h - LLVMArmNeon to LLVMIR -*- C++ -*-===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 //
-// This file implements the dialect interface for translating the LLVMArmNeon
-// dialect to LLVM IR.
+// This provides registration calls for LLVMArmNeon dialect to LLVM IR
+// translation.
 //
 //===----------------------------------------------------------------------===//
 
 #ifndef MLIR_TARGET_LLVMIR_DIALECT_LLVMARMNEON_LLVMARMNEONTOLLVMIRTRANSLATION_H
 #define MLIR_TARGET_LLVMIR_DIALECT_LLVMARMNEON_LLVMARMNEONTOLLVMIRTRANSLATION_H
 
-#include "mlir/Target/LLVMIR/LLVMTranslationInterface.h"
-
 namespace mlir {
 
-/// Implementation of the dialect interface that converts operations belonging
-/// to the LLVMArmNeon dialect to LLVM IR.
-class LLVMArmNeonDialectLLVMIRTranslationInterface
-    : public LLVMTranslationDialectInterface {
-public:
-  using LLVMTranslationDialectInterface::LLVMTranslationDialectInterface;
-
-  /// Translates the given operation to LLVM IR using the provided IR builder
-  /// and saving the state in `moduleTranslation`.
-  LogicalResult
-  convertOperation(Operation *op, llvm::IRBuilderBase &builder,
-                   LLVM::ModuleTranslation &moduleTranslation) const final;
-};
+class DialectRegistry;
+class MLIRContext;
+
+/// Register the LLVMArmNeon dialect and the translation from it to the LLVM IR
+/// in the given registry;
+void registerLLVMArmNeonDialectTranslation(DialectRegistry &registry);
+
+/// Register the LLVMArmNeon dialect and the translation from it in the registry
+/// associated with the given context.
+void registerLLVMArmNeonDialectTranslation(MLIRContext &context);
 
 } // namespace mlir
 
 #endif // MLIR_TARGET_LLVMIR_DIALECT_LLVMARMNEON_LLVMARMNEONTOLLVMIRTRANSLATION_H

mlir/include/mlir/Target/LLVMIR/Dialect/LLVMArmSVE/LLVMArmSVEToLLVMIRTranslation.h

 //===- LLVMArmSVEToLLVMIRTranslation.h - LLVMArmSVE to LLVM IR --*- C++ -*-===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 //
-// This file implements the dialect interface for translating the LLVMArmSVE
-// dialect to LLVM IR.
+// This provides registration calls for LLVMArmSVE dialect to LLVM IR
+// translation.
 //
 //===----------------------------------------------------------------------===//
 
 #ifndef MLIR_TARGET_LLVMIR_DIALECT_LLVMARMSVE_LLVMARMSVETOLLVMIRTRANSLATION_H
 #define MLIR_TARGET_LLVMIR_DIALECT_LLVMARMSVE_LLVMARMSVETOLLVMIRTRANSLATION_H
 
-#include "mlir/Target/LLVMIR/LLVMTranslationInterface.h"
-
 namespace mlir {
 
-/// Implementation of the dialect interface that converts operations belonging
-/// to the LLVMArmSVE dialect to LLVM IR.
-class LLVMArmSVEDialectLLVMIRTranslationInterface
-    : public LLVMTranslationDialectInterface {
-public:
-  using LLVMTranslationDialectInterface::LLVMTranslationDialectInterface;
-
-  /// Translates the given operation to LLVM IR using the provided IR builder
-  /// and saving the state in `moduleTranslation`.
-  LogicalResult
-  convertOperation(Operation *op, llvm::IRBuilderBase &builder,
-                   LLVM::ModuleTranslation &moduleTranslation) const final;
-};
+class DialectRegistry;
+class MLIRContext;
+
+/// Register the LLVMArmSVE dialect and the translation from it to the LLVM IR
+/// in the given registry;
+void registerLLVMArmSVEDialectTranslation(DialectRegistry &registry);
+
+/// Register the LLVMArmSVE dialect and the translation from it in the registry
+/// associated with the given context.
+void registerLLVMArmSVEDialectTranslation(MLIRContext &context);
 
 } // namespace mlir
 
 #endif // MLIR_TARGET_LLVMIR_DIALECT_LLVMARMSVE_LLVMARMSVETOLLVMIRTRANSLATION_H

mlir/include/mlir/Target/LLVMIR/Dialect/LLVMIR/LLVMToLLVMIRTranslation.h

 //===- LLVMToLLVMIRTranslation.h - LLVM Dialect to LLVM IR-------*- C++ -*-===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 //
-// This file implements the dialect interface for translating the LLVM dialect
-// to LLVM IR.
+// This provides registration calls for LLVM dialect to LLVM IR translation.
 //
 //===----------------------------------------------------------------------===//
 
 #ifndef MLIR_TARGET_LLVMIR_DIALECT_LLVMIR_LLVMTOLLVMIRTRANSLATION_H
 #define MLIR_TARGET_LLVMIR_DIALECT_LLVMIR_LLVMTOLLVMIRTRANSLATION_H
 
-#include "mlir/Target/LLVMIR/LLVMTranslationInterface.h"
-
 namespace mlir {
 
-/// Implementation of the dialect interface that converts operations belonging
-/// to the LLVM dialect to LLVM IR.
-class LLVMDialectLLVMIRTranslationInterface
-    : public LLVMTranslationDialectInterface {
-public:
-  using LLVMTranslationDialectInterface::LLVMTranslationDialectInterface;
-
-  /// Translates the given operation to LLVM IR using the provided IR builder
-  /// and saving the state in `moduleTranslation`.
-  LogicalResult
-  convertOperation(Operation *op, llvm::IRBuilderBase &builder,
-                   LLVM::ModuleTranslation &moduleTranslation) const final;
-};
+class DialectRegistry;
+class MLIRContext;
+
+/// Register the LLVM dialect and the translation from it to the LLVM IR in the
+/// given registry;
+void registerLLVMDialectTranslation(DialectRegistry &registry);
+
+/// Register the LLVM dialect and the translation from it in the registry
+/// associated with the given context.
+void registerLLVMDialectTranslation(MLIRContext &context);
 
 } // namespace mlir
 
 #endif // MLIR_TARGET_LLVMIR_DIALECT_LLVMIR_LLVMTOLLVMIRTRANSLATION_H

mlir/include/mlir/Target/LLVMIR/Dialect/NVVM/NVVMToLLVMIRTranslation.h

 //===- NVVMToLLVMIRTranslation.h - NVVM to LLVM IR --------------*- C++ -*-===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 //
-// This file implements the dialect interface for translating the NVVM
-// dialect to LLVM IR.
+// This provides registration calls for NVVM dialect to LLVM IR translation.
 //
 //===----------------------------------------------------------------------===//
 
 #ifndef MLIR_TARGET_LLVMIR_DIALECT_NVVM_NVVMTOLLVMIRTRANSLATION_H
 #define MLIR_TARGET_LLVMIR_DIALECT_NVVM_NVVMTOLLVMIRTRANSLATION_H
 
-#include "mlir/Target/LLVMIR/LLVMTranslationInterface.h"
-
 namespace mlir {
 
-/// Implementation of the dialect interface that converts operations belonging
-/// to the NVVM dialect to LLVM IR.
-class NVVMDialectLLVMIRTranslationInterface
-    : public LLVMTranslationDialectInterface {
-public:
-  using LLVMTranslationDialectInterface::LLVMTranslationDialectInterface;
-
-  /// Translates the given operation to LLVM IR using the provided IR builder
-  /// and saving the state in `moduleTranslation`.
-  LogicalResult
-  convertOperation(Operation *op, llvm::IRBuilderBase &builder,
-                   LLVM::ModuleTranslation &moduleTranslation) const final;
-
-  /// Attaches module-level metadata for functions marked as kernels.
-  LogicalResult
-  amendOperation(Operation *op, NamedAttribute attribute,
-                 LLVM::ModuleTranslation &moduleTranslation) const final;
-};
+class DialectRegistry;
+class MLIRContext;
+
+/// Register the NVVM dialect and the translation from it to the LLVM IR in the
+/// given registry;
+void registerNVVMDialectTranslation(DialectRegistry &registry);
+
+/// Register the NVVM dialect and the translation from it in the registry
+/// associated with the given context.
+void registerNVVMDialectTranslation(MLIRContext &context);
 
 } // namespace mlir
 
 #endif // MLIR_TARGET_LLVMIR_DIALECT_NVVM_NVVMTOLLVMIRTRANSLATION_H

mlir/include/mlir/Target/LLVMIR/Dialect/OpenMP/OpenMPToLLVMIRTranslation.h

 //===- OpenMPToLLVMIRTranslation.h - OpenMP Dialect to LLVM IR --*- C++ -*-===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 //
-// This file implements the dialect interface for translating the OpenMP dialect
-// to LLVM IR.
+// This provides registration calls for OpenMP dialect to LLVM IR translation.
 //
 //===----------------------------------------------------------------------===//
 
 #ifndef MLIR_TARGET_LLVMIR_DIALECT_OPENMP_OPENMPTOLLVMIRTRANSLATION_H
 #define MLIR_TARGET_LLVMIR_DIALECT_OPENMP_OPENMPTOLLVMIRTRANSLATION_H
 
-#include "mlir/Target/LLVMIR/LLVMTranslationInterface.h"
-
 namespace mlir {
 
-/// Implementation of the dialect interface that converts operations belonging
-/// to the OpenMP dialect to LLVM IR.
-class OpenMPDialectLLVMIRTranslationInterface
-    : public LLVMTranslationDialectInterface {
-public:
-  using LLVMTranslationDialectInterface::LLVMTranslationDialectInterface;
-
-  /// Translates the given operation to LLVM IR using the provided IR builder
-  /// and saving the state in `moduleTranslation`.
-  LogicalResult
-  convertOperation(Operation *op, llvm::IRBuilderBase &builder,
-                   LLVM::ModuleTranslation &moduleTranslation) const final;
-};
+class DialectRegistry;
+class MLIRContext;
+
+/// Register the OpenMP dialect and the translation from it to the LLVM IR in
+/// the given registry;
+void registerOpenMPDialectTranslation(DialectRegistry &registry);
+
+/// Register the OpenMP dialect and the translation from it in the registry
+/// associated with the given context.
+void registerOpenMPDialectTranslation(MLIRContext &context);
 
 } // namespace mlir
 
 #endif // MLIR_TARGET_LLVMIR_DIALECT_OPENMP_OPENMPTOLLVMIRTRANSLATION_H

mlir/include/mlir/Target/LLVMIR/Dialect/ROCDL/ROCDLToLLVMIRTranslation.h

 //===- ROCDLToLLVMIRTranslation.h - ROCDL to LLVM IR ------------*- C++ -*-===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 //
-// This file implements the dialect interface for translating the ROCDL
-// dialect to LLVM IR.
+// This provides registration calls for ROCDL dialect to LLVM IR translation.
 //
 //===----------------------------------------------------------------------===//
 
 #ifndef MLIR_TARGET_LLVMIR_DIALECT_ROCDL_ROCDLTOLLVMIRTRANSLATION_H
 #define MLIR_TARGET_LLVMIR_DIALECT_ROCDL_ROCDLTOLLVMIRTRANSLATION_H
 
-#include "mlir/Target/LLVMIR/LLVMTranslationInterface.h"
-
 namespace mlir {
 
-/// Implementation of the dialect interface that converts operations belonging
-/// to the ROCDL dialect to LLVM IR.
-class ROCDLDialectLLVMIRTranslationInterface
-    : public LLVMTranslationDialectInterface {
-public:
-  using LLVMTranslationDialectInterface::LLVMTranslationDialectInterface;
-
-  /// Translates the given operation to LLVM IR using the provided IR builder
-  /// and saving the state in `moduleTranslation`.
-  LogicalResult
-  convertOperation(Operation *op, llvm::IRBuilderBase &builder,
-                   LLVM::ModuleTranslation &moduleTranslation) const final;
-
-  /// Attaches module-level metadata for functions marked as kernels.
-  LogicalResult
-  amendOperation(Operation *op, NamedAttribute attribute,
-                 LLVM::ModuleTranslation &moduleTranslation) const final;
-};
+class DialectRegistry;
+class MLIRContext;
+
+/// Register the ROCDL dialect and the translation from it to the LLVM IR in the
+/// given registry;
+void registerROCDLDialectTranslation(DialectRegistry &registry);
+
+/// Register the ROCDL dialect and the translation from it in the registry
+/// associated with the given context.
+void registerROCDLDialectTranslation(MLIRContext &context);
 
 } // namespace mlir
 
 #endif // MLIR_TARGET_LLVMIR_DIALECT_ROCDL_ROCDLTOLLVMIRTRANSLATION_H

mlir/include/mlir/Target/LLVMIR/TypeTranslation.h

 
 class Type;
 class MLIRContext;
 
 namespace LLVM {
 
 namespace detail {
 class TypeToLLVMIRTranslatorImpl;
-class TypeFromLLVMIRTranslatorImpl;
 } // namespace detail
 
 /// Utility class to translate MLIR LLVM dialect types to LLVM IR. Stores the
 /// translation state, in particular any identified structure types that can be
 /// reused in further translation.
 class TypeToLLVMIRTranslator {
 public:
   TypeToLLVMIRTranslator(llvm::LLVMContext &context);
   ~TypeToLLVMIRTranslator();
 
   /// Returns the preferred alignment for the type given the data layout. Note
   /// that this will perform type conversion and store its results for future
   /// uses.
   // TODO: this should be removed when MLIR has proper data layout.
   unsigned getPreferredAlignment(Type type, const llvm::DataLayout &layout);
 
   /// Translates the given MLIR LLVM dialect type to LLVM IR.
   llvm::Type *translateType(Type type);
 
 private:
   /// Private implementation.
   std::unique_ptr<detail::TypeToLLVMIRTranslatorImpl> impl;
 };
 
-/// Utility class to translate LLVM IR types to the MLIR LLVM dialect. Stores
-/// the translation state, in particular any identified structure types that are
-/// reused across translations.
-class TypeFromLLVMIRTranslator {
-public:
-  TypeFromLLVMIRTranslator(MLIRContext &context);
-  ~TypeFromLLVMIRTranslator();
-
-  /// Translates the given LLVM IR type to the MLIR LLVM dialect.
-  Type translateType(llvm::Type *type);
-
-private:
-  /// Private implementation.
-  std::unique_ptr<detail::TypeFromLLVMIRTranslatorImpl> impl;
-};
-
 } // namespace LLVM
 } // namespace mlir
 
 #endif // MLIR_TARGET_LLVMIR_TYPETRANSLATION_H

mlir/lib/Conversion/VectorToLLVM/CMakeLists.txt

add_mlir_conversion_library(MLIRVectorToLLVM
   MLIRAVX512ToLLVM
   MLIRLLVMArmNeon
   MLIRLLVMAVX512
   MLIRArmSVE
   MLIRArmSVEToLLVM
   MLIRLLVMArmSVE
   MLIRLLVMIR
   MLIRStandardToLLVM
-  MLIRTargetLLVMIRModuleTranslation
+  MLIRTargetLLVMIRExport
   MLIRTransforms
   MLIRVector
   )

mlir/lib/ExecutionEngine/CMakeLists.txt

add_mlir_library(MLIRExecutionEngine
   Target
   Vectorize
   TransformUtils
   nativecodegen
   IPO
 
   LINK_LIBS PUBLIC
   MLIRLLVMIR
-  MLIRTargetLLVMIR
+  MLIRLLVMToLLVMIRTranslation
+  MLIRTargetLLVMIRExport
   )
 
 get_property(dialect_libs GLOBAL PROPERTY MLIR_DIALECT_LIBS)
 add_mlir_library(MLIRJitRunner
   JitRunner.cpp
 
   EXCLUDE_FROM_LIBMLIR
 
   DEPENDS
   intrinsics_gen
 
   LINK_COMPONENTS
   Core
   OrcJIT
   JITLink
 
   LINK_LIBS PUBLIC
   ${dialect_libs}
   MLIRExecutionEngine
   MLIRIR
   MLIRParser
   MLIRStandard
-  MLIRTargetLLVMIR
+  MLIRLLVMToLLVMIRTranslation
+  MLIRTargetLLVMIRExport
   MLIRTransforms
   MLIRStandardToLLVM
   MLIRSupport
 )
 
 add_mlir_library(mlir_c_runner_utils
   SHARED
   CRunnerUtils.cpp

mlir/lib/Target/CMakeLists.txt

 add_subdirectory(SPIRV)
 add_subdirectory(LLVMIR)
-
-add_mlir_translation_library(MLIRTargetLLVMIRModuleTranslation
-  LLVMIR/DebugTranslation.cpp
-  LLVMIR/ModuleTranslation.cpp
-  LLVMIR/TypeTranslation.cpp
-
-  ADDITIONAL_HEADER_DIRS
-  ${MLIR_MAIN_INCLUDE_DIR}/mlir/Target/LLVMIR
-
-  DEPENDS
-  intrinsics_gen
-
-  LINK_COMPONENTS
-  Core
-  FrontendOpenMP
-  TransformUtils
-
-  LINK_LIBS PUBLIC
-  MLIRLLVMIR
-  MLIROpenMP
-  MLIRLLVMIRTransforms
-  MLIRTranslation
-  )
-
-add_mlir_translation_library(MLIRTargetLLVMIR
-  LLVMIR/ConvertFromLLVMIR.cpp
-  LLVMIR/ConvertToLLVMIR.cpp
-
-  ADDITIONAL_HEADER_DIRS
-  ${MLIR_MAIN_INCLUDE_DIR}/mlir/Target/LLVMIR
-
-  LINK_COMPONENTS
-  Core
-  IRReader
-
-  LINK_LIBS PUBLIC
-  MLIRLLVMArmNeonToLLVMIRTranslation
-  MLIRLLVMArmSVEToLLVMIRTranslation
-  MLIRLLVMAVX512ToLLVMIRTranslation
-  MLIRLLVMToLLVMIRTranslation
-  MLIRNVVMToLLVMIRTranslation
-  MLIROpenMPToLLVMIRTranslation
-  MLIRROCDLToLLVMIRTranslation
-  MLIRTargetLLVMIRModuleTranslation
-  )

mlir/lib/Target/LLVMIR/CMakeLists.txt

 add_subdirectory(Dialect)
+
+set(LLVM_OPTIONAL_SOURCES
+  ConvertFromLLVMIR.cpp
+  ConvertToLLVMIR.cpp
+  DebugTranslation.cpp
+  ModuleTranslation.cpp
+  TypeTranslation.cpp
+  )
+
+
+add_mlir_translation_library(MLIRTargetLLVMIRExport
+  DebugTranslation.cpp
+  ModuleTranslation.cpp
+  TypeTranslation.cpp
+
+  ADDITIONAL_HEADER_DIRS
+  ${MLIR_MAIN_INCLUDE_DIR}/mlir/Target/LLVMIR
+
+  DEPENDS
+  intrinsics_gen
+
+  LINK_COMPONENTS
+  Core
+  FrontendOpenMP
+  TransformUtils
+
+  LINK_LIBS PUBLIC
+  MLIRLLVMIR
+  MLIROpenMP
+  MLIRLLVMIRTransforms
+  MLIRTranslation
+  )
+
+add_mlir_translation_library(MLIRToLLVMIRTranslationRegistration
+  ConvertToLLVMIR.cpp
+
+  LINK_LIBS PUBLIC
+  MLIRLLVMArmNeonToLLVMIRTranslation
+  MLIRLLVMArmSVEToLLVMIRTranslation
+  MLIRLLVMAVX512ToLLVMIRTranslation
+  MLIRLLVMToLLVMIRTranslation
+  MLIRNVVMToLLVMIRTranslation
+  MLIROpenMPToLLVMIRTranslation
+  MLIRROCDLToLLVMIRTranslation
+  )
+
+add_mlir_translation_library(MLIRTargetLLVMIRImport
+  ConvertFromLLVMIR.cpp
+
+  ADDITIONAL_HEADER_DIRS
+  ${MLIR_MAIN_INCLUDE_DIR}/mlir/Target/LLVMIR
+
+  LINK_COMPONENTS
+  Core
+  IRReader
+
+  LINK_LIBS PUBLIC
+  MLIRLLVMIR
+  MLIRTranslation
+  )

mlir/lib/Target/LLVMIR/ConvertFromLLVMIR.cpp

 //===----------------------------------------------------------------------===//
 
 #include "mlir/Dialect/LLVMIR/LLVMDialect.h"
 #include "mlir/IR/Builders.h"
 #include "mlir/IR/BuiltinOps.h"
 #include "mlir/IR/BuiltinTypes.h"
 #include "mlir/IR/MLIRContext.h"
 #include "mlir/Target/LLVMIR.h"
-#include "mlir/Target/LLVMIR/TypeTranslation.h"
 #include "mlir/Translation.h"
 
+#include "llvm/ADT/TypeSwitch.h"
 #include "llvm/IR/Attributes.h"
 #include "llvm/IR/Constants.h"
+#include "llvm/IR/DerivedTypes.h"
 #include "llvm/IR/Function.h"
 #include "llvm/IR/InlineAsm.h"
 #include "llvm/IR/Instructions.h"
 #include "llvm/IR/Type.h"
 #include "llvm/IRReader/IRReader.h"
 #include "llvm/Support/Error.h"
 #include "llvm/Support/SourceMgr.h"
 
 using namespace mlir;
 using namespace mlir::LLVM;
 
 #include "mlir/Dialect/LLVMIR/LLVMConversionEnumsFromLLVM.inc"
 
 // Utility to print an LLVM value as a string for passing to emitError().
 // FIXME: Diagnostic should be able to natively handle types that have
 // operator << (raw_ostream&) defined.
 static std::string diag(llvm::Value &v) {
   std::string s;
   llvm::raw_string_ostream os(s);
   os << v;
   return os.str();
 }
 
+namespace mlir {
+namespace LLVM {
+namespace detail {
+/// Support for translating LLVM IR types to MLIR LLVM dialect types.
+class TypeFromLLVMIRTranslatorImpl {
+public:
+  /// Constructs a class creating types in the given MLIR context.
+  TypeFromLLVMIRTranslatorImpl(MLIRContext &context) : context(context) {}
+
+  /// Translates the given type.
+  Type translateType(llvm::Type *type) {
+    if (knownTranslations.count(type))
+      return knownTranslations.lookup(type);
+
+    Type translated =
+        llvm::TypeSwitch<llvm::Type *, Type>(type)
+            .Case<llvm::ArrayType, llvm::FunctionType, llvm::IntegerType,
+                  llvm::PointerType, llvm::StructType, llvm::FixedVectorType,
+                  llvm::ScalableVectorType>(
+                [this](auto *type) { return this->translate(type); })
+            .Default([this](llvm::Type *type) {
+              return translatePrimitiveType(type);
+            });
+    knownTranslations.try_emplace(type, translated);
+    return translated;
+  }
+
+private:
+  /// Translates the given primitive, i.e. non-parametric in MLIR nomenclature,
+  /// type.
+  Type translatePrimitiveType(llvm::Type *type) {
+    if (type->isVoidTy())
+      return LLVM::LLVMVoidType::get(&context);
+    if (type->isHalfTy())
+      return Float16Type::get(&context);
+    if (type->isBFloatTy())
+      return BFloat16Type::get(&context);
+    if (type->isFloatTy())
+      return Float32Type::get(&context);
+    if (type->isDoubleTy())
+      return Float64Type::get(&context);
+    if (type->isFP128Ty())
+      return Float128Type::get(&context);
+    if (type->isX86_FP80Ty())
+      return Float80Type::get(&context);
+    if (type->isPPC_FP128Ty())
+      return LLVM::LLVMPPCFP128Type::get(&context);
+    if (type->isX86_MMXTy())
+      return LLVM::LLVMX86MMXType::get(&context);
+    if (type->isLabelTy())
+      return LLVM::LLVMLabelType::get(&context);
+    if (type->isMetadataTy())
+      return LLVM::LLVMMetadataType::get(&context);
+    llvm_unreachable("not a primitive type");
+  }
+
+  /// Translates the given array type.
+  Type translate(llvm::ArrayType *type) {
+    return LLVM::LLVMArrayType::get(translateType(type->getElementType()),
+                                    type->getNumElements());
+  }
+
+  /// Translates the given function type.
+  Type translate(llvm::FunctionType *type) {
+    SmallVector<Type, 8> paramTypes;
+    translateTypes(type->params(), paramTypes);
+    return LLVM::LLVMFunctionType::get(translateType(type->getReturnType()),
+                                       paramTypes, type->isVarArg());
+  }
+
+  /// Translates the given integer type.
+  Type translate(llvm::IntegerType *type) {
+    return IntegerType::get(&context, type->getBitWidth());
+  }
+
+  /// Translates the given pointer type.
+  Type translate(llvm::PointerType *type) {
+    return LLVM::LLVMPointerType::get(translateType(type->getElementType()),
+                                      type->getAddressSpace());
+  }
+
+  /// Translates the given structure type.
+  Type translate(llvm::StructType *type) {
+    SmallVector<Type, 8> subtypes;
+    if (type->isLiteral()) {
+      translateTypes(type->subtypes(), subtypes);
+      return LLVM::LLVMStructType::getLiteral(&context, subtypes,
+                                              type->isPacked());
+    }
+
+    if (type->isOpaque())
+      return LLVM::LLVMStructType::getOpaque(type->getName(), &context);
+
+    LLVM::LLVMStructType translated =
+        LLVM::LLVMStructType::getIdentified(&context, type->getName());
+    knownTranslations.try_emplace(type, translated);
+    translateTypes(type->subtypes(), subtypes);
+    LogicalResult bodySet = translated.setBody(subtypes, type->isPacked());
+    assert(succeeded(bodySet) &&
+           "could not set the body of an identified struct");
+    (void)bodySet;
+    return translated;
+  }
+
+  /// Translates the given fixed-vector type.
+  Type translate(llvm::FixedVectorType *type) {
+    return LLVM::getFixedVectorType(translateType(type->getElementType()),
+                                    type->getNumElements());
+  }
+
+  /// Translates the given scalable-vector type.
+  Type translate(llvm::ScalableVectorType *type) {
+    return LLVM::LLVMScalableVectorType::get(
+        translateType(type->getElementType()), type->getMinNumElements());
+  }
+
+  /// Translates a list of types.
+  void translateTypes(ArrayRef<llvm::Type *> types,
+                      SmallVectorImpl<Type> &result) {
+    result.reserve(result.size() + types.size());
+    for (llvm::Type *type : types)
+      result.push_back(translateType(type));
+  }
+
+  /// Map of known translations. Serves as a cache and as recursion stopper for
+  /// translating recursive structs.
+  llvm::DenseMap<llvm::Type *, Type> knownTranslations;
+
+  /// The context in which MLIR types are created.
+  MLIRContext &context;
+};
+} // end namespace detail
+
+/// Utility class to translate LLVM IR types to the MLIR LLVM dialect. Stores
+/// the translation state, in particular any identified structure types that are
+/// reused across translations.
+class TypeFromLLVMIRTranslator {
+public:
+  TypeFromLLVMIRTranslator(MLIRContext &context);
+  ~TypeFromLLVMIRTranslator();
+
+  /// Translates the given LLVM IR type to the MLIR LLVM dialect.
+  Type translateType(llvm::Type *type);
+
+private:
+  /// Private implementation.
+  std::unique_ptr<detail::TypeFromLLVMIRTranslatorImpl> impl;
+};
+
+} // end namespace LLVM
+} // end namespace mlir
+
+LLVM::TypeFromLLVMIRTranslator::TypeFromLLVMIRTranslator(MLIRContext &context)
+    : impl(new detail::TypeFromLLVMIRTranslatorImpl(context)) {}
+
+LLVM::TypeFromLLVMIRTranslator::~TypeFromLLVMIRTranslator() {}
+
+Type LLVM::TypeFromLLVMIRTranslator::translateType(llvm::Type *type) {
+  return impl->translateType(type);
+}
+
 // Handles importing globals and functions from an LLVM module.
 namespace {
 class Importer {
 public:
   Importer(MLIRContext *context, ModuleOp module)
       : b(context), context(context), module(module),
         unknownLoc(FileLineColLoc::get("imported-bitcode", 0, 0, context)),
         typeTranslator(*context) {

mlir/lib/Target/LLVMIR/ConvertToLLVMIR.cpp

 //===- ConvertToLLVMIR.cpp - MLIR to LLVM IR conversion -------------------===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 //
 // This file implements a translation between the MLIR LLVM dialect and LLVM IR.
 //
 //===----------------------------------------------------------------------===//
 
-#include "mlir/Target/LLVMIR.h"
-
-#include "mlir/Dialect/LLVMIR/LLVMAVX512Dialect.h"
-#include "mlir/Dialect/LLVMIR/LLVMArmNeonDialect.h"
-#include "mlir/Dialect/LLVMIR/LLVMArmSVEDialect.h"
-#include "mlir/Dialect/LLVMIR/NVVMDialect.h"
-#include "mlir/Dialect/LLVMIR/ROCDLDialect.h"
-#include "mlir/Dialect/OpenMP/OpenMPDialect.h"
-#include "mlir/Target/LLVMIR/Dialect/LLVMAVX512/LLVMAVX512ToLLVMIRTranslation.h"
-#include "mlir/Target/LLVMIR/Dialect/LLVMArmNeon/LLVMArmNeonToLLVMIRTranslation.h"
-#include "mlir/Target/LLVMIR/Dialect/LLVMArmSVE/LLVMArmSVEToLLVMIRTranslation.h"
-#include "mlir/Target/LLVMIR/Dialect/LLVMIR/LLVMToLLVMIRTranslation.h"
-#include "mlir/Target/LLVMIR/Dialect/NVVM/NVVMToLLVMIRTranslation.h"
-#include "mlir/Target/LLVMIR/Dialect/OpenMP/OpenMPToLLVMIRTranslation.h"
-#include "mlir/Target/LLVMIR/Dialect/ROCDL/ROCDLToLLVMIRTranslation.h"
-#include "mlir/Target/LLVMIR/ModuleTranslation.h"
+#include "mlir/IR/BuiltinOps.h"
+#include "mlir/Target/LLVMIR/Dialect/All.h"
+#include "mlir/Target/LLVMIR/Export.h"
 #include "mlir/Translation.h"
-
-#include "llvm/ADT/StringRef.h"
+#include "llvm/IR/LLVMContext.h"
 #include "llvm/IR/Module.h"
-#include "llvm/IR/Verifier.h"
-#include "llvm/Support/ToolOutputFile.h"
 
 using namespace mlir;
 
-void mlir::registerLLVMDialectTranslation(DialectRegistry &registry) {
-  registry.insert<LLVM::LLVMDialect>();
-  registry.addDialectInterface<LLVM::LLVMDialect,
-                               LLVMDialectLLVMIRTranslationInterface>();
-}
-
-void mlir::registerLLVMDialectTranslation(MLIRContext &context) {
-  auto *dialect = context.getLoadedDialect<LLVM::LLVMDialect>();
-  if (!dialect || dialect->getRegisteredInterface<
-                      LLVMDialectLLVMIRTranslationInterface>() == nullptr) {
-    DialectRegistry registry;
-    registry.insert<LLVM::LLVMDialect>();
-    registry.addDialectInterface<LLVM::LLVMDialect,
-                                 LLVMDialectLLVMIRTranslationInterface>();
-    context.appendDialectRegistry(registry);
-  }
-}
-
 namespace mlir {
 void registerToLLVMIRTranslation() {
   TranslateFromMLIRRegistration registration(
       "mlir-to-llvmir",
       [](ModuleOp module, raw_ostream &output) {
         llvm::LLVMContext llvmContext;
         auto llvmModule = translateModuleToLLVMIR(module, llvmContext);
         if (!llvmModule)
           return failure();
 
         llvmModule->print(output, nullptr);
         return success();
       },
       [](DialectRegistry &registry) {
-        registry.insert<omp::OpenMPDialect, LLVM::LLVMAVX512Dialect,
-                        LLVM::LLVMArmSVEDialect, LLVM::LLVMArmNeonDialect,
-                        NVVM::NVVMDialect, ROCDL::ROCDLDialect>();
-        registry.addDialectInterface<omp::OpenMPDialect,
-                                     OpenMPDialectLLVMIRTranslationInterface>();
-        registry
-            .addDialectInterface<LLVM::LLVMAVX512Dialect,
-                                 LLVMAVX512DialectLLVMIRTranslationInterface>();
-        registry.addDialectInterface<
-            LLVM::LLVMArmNeonDialect,
-            LLVMArmNeonDialectLLVMIRTranslationInterface>();
-        registry
-            .addDialectInterface<LLVM::LLVMArmSVEDialect,
-                                 LLVMArmSVEDialectLLVMIRTranslationInterface>();
-        registry.addDialectInterface<NVVM::NVVMDialect,
-                                     NVVMDialectLLVMIRTranslationInterface>();
-        registry.addDialectInterface<ROCDL::ROCDLDialect,
-                                     ROCDLDialectLLVMIRTranslationInterface>();
-        registerLLVMDialectTranslation(registry);
+        registerAllToLLVMIRTranslations(registry);
       });
 }
 } // namespace mlir

mlir/lib/Target/LLVMIR/Dialect/LLVMAVX512/CMakeLists.txt

 add_mlir_translation_library(MLIRLLVMAVX512ToLLVMIRTranslation
   LLVMAVX512ToLLVMIRTranslation.cpp
 
   DEPENDS
   MLIRLLVMAVX512ConversionsIncGen
 
   LINK_COMPONENTS
   Core
 
   LINK_LIBS PUBLIC
   MLIRIR
   MLIRLLVMAVX512
   MLIRLLVMIR
   MLIRSupport
-  MLIRTargetLLVMIRModuleTranslation
+  MLIRTargetLLVMIRExport
   )

mlir/lib/Target/LLVMIR/Dialect/LLVMAVX512/LLVMAVX512ToLLVMIRTranslation.cpp

 #include "mlir/Target/LLVMIR/ModuleTranslation.h"
 
 #include "llvm/IR/IRBuilder.h"
 #include "llvm/IR/IntrinsicsX86.h"
 
 using namespace mlir;
 using namespace mlir::LLVM;
 
+namespace {
+/// Implementation of the dialect interface that converts operations belonging
+/// to the LLVMAVX512 dialect to LLVM IR.
+class LLVMAVX512DialectLLVMIRTranslationInterface
+    : public LLVMTranslationDialectInterface {
+public:
+  using LLVMTranslationDialectInterface::LLVMTranslationDialectInterface;
+
+  /// Translates the given operation to LLVM IR using the provided IR builder
+  /// and saving the state in `moduleTranslation`.
   LogicalResult
-mlir::LLVMAVX512DialectLLVMIRTranslationInterface::convertOperation(
-    Operation *op, llvm::IRBuilderBase &builder,
-    LLVM::ModuleTranslation &moduleTranslation) const {
+  convertOperation(Operation *op, llvm::IRBuilderBase &builder,
+                   LLVM::ModuleTranslation &moduleTranslation) const final {
     Operation &opInst = *op;
 #include "mlir/Dialect/LLVMIR/LLVMAVX512Conversions.inc"
 
     return failure();
   }
+};
+} // end namespace
+
+void mlir::registerLLVMAVX512DialectTranslation(DialectRegistry &registry) {
+  registry.insert<LLVM::LLVMAVX512Dialect>();
+  registry.addDialectInterface<LLVM::LLVMAVX512Dialect,
+                               LLVMAVX512DialectLLVMIRTranslationInterface>();
+}
+
+void mlir::registerLLVMAVX512DialectTranslation(MLIRContext &context) {
+  DialectRegistry registry;
+  registerLLVMAVX512DialectTranslation(registry);
+  context.appendDialectRegistry(registry);
+}

mlir/lib/Target/LLVMIR/Dialect/LLVMArmNeon/CMakeLists.txt

 add_mlir_translation_library(MLIRLLVMArmNeonToLLVMIRTranslation
   LLVMArmNeonToLLVMIRTranslation.cpp
 
   DEPENDS
   MLIRLLVMArmNeonConversionsIncGen
 
   LINK_COMPONENTS
   Core
 
   LINK_LIBS PUBLIC
   MLIRIR
   MLIRLLVMArmNeon
   MLIRLLVMIR
   MLIRSupport
-  MLIRTargetLLVMIRModuleTranslation
+  MLIRTargetLLVMIRExport
   )

mlir/lib/Target/LLVMIR/Dialect/LLVMArmNeon/LLVMArmNeonToLLVMIRTranslation.cpp

 #include "mlir/Target/LLVMIR/ModuleTranslation.h"
 
 #include "llvm/IR/IRBuilder.h"
 #include "llvm/IR/IntrinsicsAArch64.h"
 
 using namespace mlir;
 using namespace mlir::LLVM;
 
+namespace {
+/// Implementation of the dialect interface that converts operations belonging
+/// to the LLVMArmNeon dialect to LLVM IR.
+class LLVMArmNeonDialectLLVMIRTranslationInterface
+    : public LLVMTranslationDialectInterface {
+public:
+  using LLVMTranslationDialectInterface::LLVMTranslationDialectInterface;
+
+  /// Translates the given operation to LLVM IR using the provided IR builder
+  /// and saving the state in `moduleTranslation`.
   LogicalResult
-mlir::LLVMArmNeonDialectLLVMIRTranslationInterface::convertOperation(
-    Operation *op, llvm::IRBuilderBase &builder,
-    LLVM::ModuleTranslation &moduleTranslation) const {
+  convertOperation(Operation *op, llvm::IRBuilderBase &builder,
+                   LLVM::ModuleTranslation &moduleTranslation) const final {
     Operation &opInst = *op;
 #include "mlir/Dialect/LLVMIR/LLVMArmNeonConversions.inc"
 
     return failure();
   }
+};
+} // end namespace
+
+void mlir::registerLLVMArmNeonDialectTranslation(DialectRegistry &registry) {
+  registry.insert<LLVM::LLVMArmNeonDialect>();
+  registry.addDialectInterface<LLVM::LLVMArmNeonDialect,
+                               LLVMArmNeonDialectLLVMIRTranslationInterface>();
+}
+
+void mlir::registerLLVMArmNeonDialectTranslation(MLIRContext &context) {
+  DialectRegistry registry;
+  registerLLVMArmNeonDialectTranslation(registry);
+  context.appendDialectRegistry(registry);
+}

mlir/lib/Target/LLVMIR/Dialect/LLVMArmSVE/CMakeLists.txt

 add_mlir_translation_library(MLIRLLVMArmSVEToLLVMIRTranslation
   LLVMArmSVEToLLVMIRTranslation.cpp
 
   DEPENDS
   MLIRLLVMArmSVEConversionsIncGen
 
   LINK_COMPONENTS
   Core
 
   LINK_LIBS PUBLIC
   MLIRIR
   MLIRLLVMArmSVE
   MLIRLLVMIR
   MLIRSupport
-  MLIRTargetLLVMIRModuleTranslation
+  MLIRTargetLLVMIRExport
   )

mlir/lib/Target/LLVMIR/Dialect/LLVMArmSVE/LLVMArmSVEToLLVMIRTranslation.cpp

 #include "mlir/Target/LLVMIR/ModuleTranslation.h"
 
 #include "llvm/IR/IRBuilder.h"
 #include "llvm/IR/IntrinsicsAArch64.h"
 
 using namespace mlir;
 using namespace mlir::LLVM;
 
+namespace {
+/// Implementation of the dialect interface that converts operations belonging
+/// to the LLVMArmSVE dialect to LLVM IR.
+class LLVMArmSVEDialectLLVMIRTranslationInterface
+    : public LLVMTranslationDialectInterface {
+public:
+  using LLVMTranslationDialectInterface::LLVMTranslationDialectInterface;
+
+  /// Translates the given operation to LLVM IR using the provided IR builder
+  /// and saving the state in `moduleTranslation`.
   LogicalResult
-mlir::LLVMArmSVEDialectLLVMIRTranslationInterface::convertOperation(
-    Operation *op, llvm::IRBuilderBase &builder,
-    LLVM::ModuleTranslation &moduleTranslation) const {
+  convertOperation(Operation *op, llvm::IRBuilderBase &builder,
+                   LLVM::ModuleTranslation &moduleTranslation) const final {
     Operation &opInst = *op;
 #include "mlir/Dialect/LLVMIR/LLVMArmSVEConversions.inc"
 
     return failure();
   }
+};
+} // end namespace
+
+void mlir::registerLLVMArmSVEDialectTranslation(DialectRegistry &registry) {
+  registry.insert<LLVM::LLVMArmSVEDialect>();
+  registry.addDialectInterface<LLVM::LLVMArmSVEDialect,
+                               LLVMArmSVEDialectLLVMIRTranslationInterface>();
+}
+
+void mlir::registerLLVMArmSVEDialectTranslation(MLIRContext &context) {
+  DialectRegistry registry;
+  registerLLVMArmSVEDialectTranslation(registry);
+  context.appendDialectRegistry(registry);
+}

mlir/lib/Target/LLVMIR/Dialect/LLVMIR/CMakeLists.txt

 add_mlir_translation_library(MLIRLLVMToLLVMIRTranslation
   LLVMToLLVMIRTranslation.cpp
 
   LINK_COMPONENTS
   Core
 
   LINK_LIBS PUBLIC
   MLIRIR
   MLIRLLVMIR
   MLIRSupport
-  MLIRTargetLLVMIRModuleTranslation
+  MLIRTargetLLVMIRExport
   )

mlir/lib/Target/LLVMIR/Dialect/LLVMIR/LLVMToLLVMIRTranslation.cpp

moduleTranslation.mapValue(
         global ? moduleTranslation.lookupGlobal(global)
                : moduleTranslation.lookupFunction(function.getName()));
     return success();
   }
 
   return failure();
 }
 
-LogicalResult mlir::LLVMDialectLLVMIRTranslationInterface::convertOperation(
-    Operation *op, llvm::IRBuilderBase &builder,
-    LLVM::ModuleTranslation &moduleTranslation) const {
+namespace {
+/// Implementation of the dialect interface that converts operations belonging
+/// to the LLVM dialect to LLVM IR.
+class LLVMDialectLLVMIRTranslationInterface
+    : public LLVMTranslationDialectInterface {
+public:
+  using LLVMTranslationDialectInterface::LLVMTranslationDialectInterface;
+
+  /// Translates the given operation to LLVM IR using the provided IR builder
+  /// and saving the state in `moduleTranslation`.
+  LogicalResult
+  convertOperation(Operation *op, llvm::IRBuilderBase &builder,
+                   LLVM::ModuleTranslation &moduleTranslation) const final {
     return convertOperationImpl(*op, builder, moduleTranslation);
   }
+};
+} // end namespace
+
+void mlir::registerLLVMDialectTranslation(DialectRegistry &registry) {
+  registry.insert<LLVM::LLVMDialect>();
+  registry.addDialectInterface<LLVM::LLVMDialect,
+                               LLVMDialectLLVMIRTranslationInterface>();
+}
+
+void mlir::registerLLVMDialectTranslation(MLIRContext &context) {
+  DialectRegistry registry;
+  registerLLVMDialectTranslation(registry);
+  context.appendDialectRegistry(registry);
+}

mlir/lib/Target/LLVMIR/Dialect/NVVM/CMakeLists.txt

 add_mlir_translation_library(MLIRNVVMToLLVMIRTranslation
   NVVMToLLVMIRTranslation.cpp
 
   DEPENDS
   MLIRNVVMConversionsIncGen
 
   LINK_COMPONENTS
   Core
 
   LINK_LIBS PUBLIC
   MLIRIR
   MLIRLLVMIR
   MLIRNVVMIR
   MLIRSupport
-  MLIRTargetLLVMIRModuleTranslation
+  MLIRTargetLLVMIRExport
   )

mlir/lib/Target/LLVMIR/Dialect/NVVM/NVVMToLLVMIRTranslation.cpp

if (withPredicate) {
     resultType = cast<llvm::StructType>(resultType)->getElementType(0);
     return resultType->isFloatTy() ? llvm::Intrinsic::nvvm_shfl_sync_bfly_f32p
                                    : llvm::Intrinsic::nvvm_shfl_sync_bfly_i32p;
   }
   return resultType->isFloatTy() ? llvm::Intrinsic::nvvm_shfl_sync_bfly_f32
                                  : llvm::Intrinsic::nvvm_shfl_sync_bfly_i32;
 }
 
-LogicalResult mlir::NVVMDialectLLVMIRTranslationInterface::convertOperation(
-    Operation *op, llvm::IRBuilderBase &builder,
-    LLVM::ModuleTranslation &moduleTranslation) const {
+namespace {
+/// Implementation of the dialect interface that converts operations belonging
+/// to the NVVM dialect to LLVM IR.
+class NVVMDialectLLVMIRTranslationInterface
+    : public LLVMTranslationDialectInterface {
+public:
+  using LLVMTranslationDialectInterface::LLVMTranslationDialectInterface;
+
+  /// Translates the given operation to LLVM IR using the provided IR builder
+  /// and saving the state in `moduleTranslation`.
+  LogicalResult
+  convertOperation(Operation *op, llvm::IRBuilderBase &builder,
+                   LLVM::ModuleTranslation &moduleTranslation) const final {
     Operation &opInst = *op;
 #include "mlir/Dialect/LLVMIR/NVVMConversions.inc"
 
     return failure();
   }
 
-LogicalResult mlir::NVVMDialectLLVMIRTranslationInterface::amendOperation(
-    Operation *op, NamedAttribute attribute,
-    LLVM::ModuleTranslation &moduleTranslation) const {
+  /// Attaches module-level metadata for functions marked as kernels.
+  LogicalResult
+  amendOperation(Operation *op, NamedAttribute attribute,
+                 LLVM::ModuleTranslation &moduleTranslation) const final {
     if (attribute.first == NVVM::NVVMDialect::getKernelFuncAttrName()) {
-    auto func = cast<LLVM::LLVMFuncOp>(op);
+      auto func = dyn_cast<LLVM::LLVMFuncOp>(op);
+      if (!func)
+        return failure();
+
       llvm::LLVMContext &llvmContext = moduleTranslation.getLLVMContext();
-    llvm::Function *llvmFunc = moduleTranslation.lookupFunction(func.getName());
+      llvm::Function *llvmFunc =
+          moduleTranslation.lookupFunction(func.getName());
       llvm::Metadata *llvmMetadata[] = {
           llvm::ValueAsMetadata::get(llvmFunc),
           llvm::MDString::get(llvmContext, "kernel"),
           llvm::ValueAsMetadata::get(
               llvm::ConstantInt::get(llvm::Type::getInt32Ty(llvmContext), 1))};
       llvm::MDNode *llvmMetadataNode =
           llvm::MDNode::get(llvmContext, llvmMetadata);
       moduleTranslation.getOrInsertNamedModuleMetadata("nvvm.annotations")
           ->addOperand(llvmMetadataNode);
     }
     return success();
   }
+};
+} // end namespace
+
+void mlir::registerNVVMDialectTranslation(DialectRegistry &registry) {
+  registry.insert<NVVM::NVVMDialect>();
+  registry.addDialectInterface<NVVM::NVVMDialect,
+                               NVVMDialectLLVMIRTranslationInterface>();
+}
+
+void mlir::registerNVVMDialectTranslation(MLIRContext &context) {
+  DialectRegistry registry;
+  registerNVVMDialectTranslation(registry);
+  context.appendDialectRegistry(registry);
+}

mlir/lib/Target/LLVMIR/Dialect/OpenMP/CMakeLists.txt

 add_mlir_translation_library(MLIROpenMPToLLVMIRTranslation
   OpenMPToLLVMIRTranslation.cpp
 
   LINK_COMPONENTS
   Core
 
   LINK_LIBS PUBLIC
   MLIRIR
   MLIRLLVMIR
   MLIROpenMP
   MLIRSupport
-  MLIRTargetLLVMIRModuleTranslation
+  MLIRTargetLLVMIRExport
   )

mlir/lib/Target/LLVMIR/Dialect/OpenMP/OpenMPToLLVMIRTranslation.cpp

convertOmpMaster(Operation &opInst, llvm::IRBuilderBase &builder,
   llvm::OpenMPIRBuilder::LocationDescription ompLoc(
       builder.saveIP(), builder.getCurrentDebugLocation());
   builder.restoreIP(moduleTranslation.getOpenMPBuilder()->createMaster(
       ompLoc, bodyGenCB, finiCB));
   return success();
 }
 
 /// Converts an OpenMP workshare loop into LLVM IR using OpenMPIRBuilder.
-LogicalResult convertOmpWsLoop(Operation &opInst, llvm::IRBuilderBase &builder,
-                               LLVM::ModuleTranslation &moduleTranslation) {
+static LogicalResult
+convertOmpWsLoop(Operation &opInst, llvm::IRBuilderBase &builder,
   auto loop = cast<omp::WsLoopOp>(opInst);
   // TODO: this should be in the op verifier instead.
   if (loop.lowerBound().empty())
     return failure();
 
   if (loop.getNumLoops() != 1)
     return opInst.emitOpError("collapsed loops not yet supported");
 
convertOmpWsLoop(Operation &opInst, llvm::IRBuilderBase &builder,
   loopInfo = moduleTranslation.getOpenMPBuilder()->createStaticWorkshareLoop(
       ompLoc, loopInfo, allocaIP, !loop.nowait(), chunk);
 
   // Continue building IR after the loop.
   builder.restoreIP(loopInfo->getAfterIP());
   return success();
 }
 
+namespace {
+
+/// Implementation of the dialect interface that converts operations belonging
+/// to the OpenMP dialect to LLVM IR.
+class OpenMPDialectLLVMIRTranslationInterface
+    : public LLVMTranslationDialectInterface {
+public:
+  using LLVMTranslationDialectInterface::LLVMTranslationDialectInterface;
+
+  /// Translates the given operation to LLVM IR using the provided IR builder
+  /// and saving the state in `moduleTranslation`.
+  LogicalResult
+  convertOperation(Operation *op, llvm::IRBuilderBase &builder,
+                   LLVM::ModuleTranslation &moduleTranslation) const final;
+};
+
+} // end namespace
+
 /// Given an OpenMP MLIR operation, create the corresponding LLVM IR
 /// (including OpenMP runtime calls).
-LogicalResult mlir::OpenMPDialectLLVMIRTranslationInterface::convertOperation(
+LogicalResult OpenMPDialectLLVMIRTranslationInterface::convertOperation(
     Operation *op, llvm::IRBuilderBase &builder,
     LLVM::ModuleTranslation &moduleTranslation) const {
 
   llvm::OpenMPIRBuilder *ompBuilder = moduleTranslation.getOpenMPBuilder();
 
   return llvm::TypeSwitch<Operation *, LogicalResult>(op)
       .Case([&](omp::BarrierOp) {
         ompBuilder->createBarrier(builder.saveIP(), llvm::omp::OMPD_barrier);
return llvm::TypeSwitch<Operation *, LogicalResult>(op)
                "unexpected OpenMP terminator with operands");
         return success();
       })
       .Default([&](Operation *inst) {
         return inst->emitError("unsupported OpenMP operation: ")
                << inst->getName();
       });
 }
+
+void mlir::registerOpenMPDialectTranslation(DialectRegistry &registry) {
+  registry.insert<omp::OpenMPDialect>();
+  registry.addDialectInterface<omp::OpenMPDialect,
+                               OpenMPDialectLLVMIRTranslationInterface>();
+}
+
+void mlir::registerOpenMPDialectTranslation(MLIRContext &context) {
+  DialectRegistry registry;
+  registerOpenMPDialectTranslation(registry);
+  context.appendDialectRegistry(registry);
+}

mlir/lib/Target/LLVMIR/Dialect/ROCDL/CMakeLists.txt

 add_mlir_translation_library(MLIRROCDLToLLVMIRTranslation
   ROCDLToLLVMIRTranslation.cpp
 
   DEPENDS
   MLIRROCDLConversionsIncGen
 
   LINK_COMPONENTS
   Core
 
   LINK_LIBS PUBLIC
   MLIRIR
   MLIRLLVMIR
   MLIRROCDLIR
   MLIRSupport
-  MLIRTargetLLVMIRModuleTranslation
+  MLIRTargetLLVMIRExport
   )

mlir/lib/Target/LLVMIR/Dialect/ROCDL/ROCDLToLLVMIRTranslation.cpp

llvm::FunctionType *function_type = llvm::FunctionType::get(
       false);                                       // no variadic arguments.
   llvm::Function *fn = dyn_cast<llvm::Function>(
       module->getOrInsertFunction(fn_name, function_type).getCallee());
   llvm::Value *fn_op0 = llvm::ConstantInt::get(
       llvm::Type::getInt32Ty(module->getContext()), parameter);
   return builder.CreateCall(fn, ArrayRef<llvm::Value *>(fn_op0));
 }
 
-LogicalResult mlir::ROCDLDialectLLVMIRTranslationInterface::convertOperation(
-    Operation *op, llvm::IRBuilderBase &builder,
-    LLVM::ModuleTranslation &moduleTranslation) const {
+namespace {
+/// Implementation of the dialect interface that converts operations belonging
+/// to the ROCDL dialect to LLVM IR.
+class ROCDLDialectLLVMIRTranslationInterface
+    : public LLVMTranslationDialectInterface {
+public:
+  using LLVMTranslationDialectInterface::LLVMTranslationDialectInterface;
+
+  /// Translates the given operation to LLVM IR using the provided IR builder
+  /// and saving the state in `moduleTranslation`.
+  LogicalResult
+  convertOperation(Operation *op, llvm::IRBuilderBase &builder,
+                   LLVM::ModuleTranslation &moduleTranslation) const final {
     Operation &opInst = *op;
 #include "mlir/Dialect/LLVMIR/ROCDLConversions.inc"
 
     return failure();
   }
 
-LogicalResult mlir::ROCDLDialectLLVMIRTranslationInterface::amendOperation(
-    Operation *op, NamedAttribute attribute,
-    LLVM::ModuleTranslation &moduleTranslation) const {
+  /// Attaches module-level metadata for functions marked as kernels.
+  LogicalResult
+  amendOperation(Operation *op, NamedAttribute attribute,
+                 LLVM::ModuleTranslation &moduleTranslation) const final {
     if (attribute.first == ROCDL::ROCDLDialect::getKernelFuncAttrName()) {
-    auto func = cast<LLVM::LLVMFuncOp>(op);
+      auto func = dyn_cast<LLVM::LLVMFuncOp>(op);
+      if (!func)
+        return failure();
 
       // For GPU kernels,
       // 1. Insert AMDGPU_KERNEL calling convention.
       // 2. Insert amdgpu-flat-workgroup-size(1, 1024) attribute.
-    llvm::Function *llvmFunc = moduleTranslation.lookupFunction(func.getName());
+      llvm::Function *llvmFunc =
+          moduleTranslation.lookupFunction(func.getName());
       llvmFunc->setCallingConv(llvm::CallingConv::AMDGPU_KERNEL);
       llvmFunc->addFnAttr("amdgpu-flat-work-group-size", "1, 1024");
     }
     return success();
   }
+};
+} // end namespace
+
+void mlir::registerROCDLDialectTranslation(DialectRegistry &registry) {
+  registry.insert<ROCDL::ROCDLDialect>();
+  registry.addDialectInterface<ROCDL::ROCDLDialect,
+                               ROCDLDialectLLVMIRTranslationInterface>();
+}
+
+void mlir::registerROCDLDialectTranslation(MLIRContext &context) {
+  DialectRegistry registry;
+  registerROCDLDialectTranslation(registry);
+  context.appendDialectRegistry(registry);
+}

mlir/lib/Target/LLVMIR/TypeTranslation.cpp

 llvm::Type *LLVM::TypeToLLVMIRTranslator::translateType(Type type) {
   return impl->translateType(type);
 }
 
 unsigned LLVM::TypeToLLVMIRTranslator::getPreferredAlignment(
     Type type, const llvm::DataLayout &layout) {
   return layout.getPrefTypeAlignment(translateType(type));
 }
-
-namespace mlir {
-namespace LLVM {
-namespace detail {
-/// Support for translating LLVM IR types to MLIR LLVM dialect types.
-class TypeFromLLVMIRTranslatorImpl {
-public:
-  /// Constructs a class creating types in the given MLIR context.
-  TypeFromLLVMIRTranslatorImpl(MLIRContext &context) : context(context) {}
-
-  /// Translates the given type.
-  Type translateType(llvm::Type *type) {
-    if (knownTranslations.count(type))
-      return knownTranslations.lookup(type);
-
-    Type translated =
-        llvm::TypeSwitch<llvm::Type *, Type>(type)
-            .Case<llvm::ArrayType, llvm::FunctionType, llvm::IntegerType,
-                  llvm::PointerType, llvm::StructType, llvm::FixedVectorType,
-                  llvm::ScalableVectorType>(
-                [this](auto *type) { return this->translate(type); })
-            .Default([this](llvm::Type *type) {
-              return translatePrimitiveType(type);
-            });
-    knownTranslations.try_emplace(type, translated);
-    return translated;
-  }
-
-private:
-  /// Translates the given primitive, i.e. non-parametric in MLIR nomenclature,
-  /// type.
-  Type translatePrimitiveType(llvm::Type *type) {
-    if (type->isVoidTy())
-      return LLVM::LLVMVoidType::get(&context);
-    if (type->isHalfTy())
-      return Float16Type::get(&context);
-    if (type->isBFloatTy())
-      return BFloat16Type::get(&context);
-    if (type->isFloatTy())
-      return Float32Type::get(&context);
-    if (type->isDoubleTy())
-      return Float64Type::get(&context);
-    if (type->isFP128Ty())
-      return Float128Type::get(&context);
-    if (type->isX86_FP80Ty())
-      return Float80Type::get(&context);
-    if (type->isPPC_FP128Ty())
-      return LLVM::LLVMPPCFP128Type::get(&context);
-    if (type->isX86_MMXTy())
-      return LLVM::LLVMX86MMXType::get(&context);
-    if (type->isLabelTy())
-      return LLVM::LLVMLabelType::get(&context);
-    if (type->isMetadataTy())
-      return LLVM::LLVMMetadataType::get(&context);
-    llvm_unreachable("not a primitive type");
-  }
-
-  /// Translates the given array type.
-  Type translate(llvm::ArrayType *type) {
-    return LLVM::LLVMArrayType::get(translateType(type->getElementType()),
-                                    type->getNumElements());
-  }
-
-  /// Translates the given function type.
-  Type translate(llvm::FunctionType *type) {
-    SmallVector<Type, 8> paramTypes;
-    translateTypes(type->params(), paramTypes);
-    return LLVM::LLVMFunctionType::get(translateType(type->getReturnType()),
-                                       paramTypes, type->isVarArg());
-  }
-
-  /// Translates the given integer type.
-  Type translate(llvm::IntegerType *type) {
-    return IntegerType::get(&context, type->getBitWidth());
-  }
-
-  /// Translates the given pointer type.
-  Type translate(llvm::PointerType *type) {
-    return LLVM::LLVMPointerType::get(translateType(type->getElementType()),
-                                      type->getAddressSpace());
-  }
-
-  /// Translates the given structure type.
-  Type translate(llvm::StructType *type) {
-    SmallVector<Type, 8> subtypes;
-    if (type->isLiteral()) {
-      translateTypes(type->subtypes(), subtypes);
-      return LLVM::LLVMStructType::getLiteral(&context, subtypes,
-                                              type->isPacked());
-    }
-
-    if (type->isOpaque())
-      return LLVM::LLVMStructType::getOpaque(type->getName(), &context);
-
-    LLVM::LLVMStructType translated =
-        LLVM::LLVMStructType::getIdentified(&context, type->getName());
-    knownTranslations.try_emplace(type, translated);
-    translateTypes(type->subtypes(), subtypes);
-    LogicalResult bodySet = translated.setBody(subtypes, type->isPacked());
-    assert(succeeded(bodySet) &&
-           "could not set the body of an identified struct");
-    (void)bodySet;
-    return translated;
-  }
-
-  /// Translates the given fixed-vector type.
-  Type translate(llvm::FixedVectorType *type) {
-    return LLVM::getFixedVectorType(translateType(type->getElementType()),
-                                    type->getNumElements());
-  }
-
-  /// Translates the given scalable-vector type.
-  Type translate(llvm::ScalableVectorType *type) {
-    return LLVM::LLVMScalableVectorType::get(
-        translateType(type->getElementType()), type->getMinNumElements());
-  }
-
-  /// Translates a list of types.
-  void translateTypes(ArrayRef<llvm::Type *> types,
-                      SmallVectorImpl<Type> &result) {
-    result.reserve(result.size() + types.size());
-    for (llvm::Type *type : types)
-      result.push_back(translateType(type));
-  }
-
-  /// Map of known translations. Serves as a cache and as recursion stopper for
-  /// translating recursive structs.
-  llvm::DenseMap<llvm::Type *, Type> knownTranslations;
-
-  /// The context in which MLIR types are created.
-  MLIRContext &context;
-};
-} // end namespace detail
-} // end namespace LLVM
-} // end namespace mlir
-
-LLVM::TypeFromLLVMIRTranslator::TypeFromLLVMIRTranslator(MLIRContext &context)
-    : impl(new detail::TypeFromLLVMIRTranslatorImpl(context)) {}
-
-LLVM::TypeFromLLVMIRTranslator::~TypeFromLLVMIRTranslator() {}
-
-Type LLVM::TypeFromLLVMIRTranslator::translateType(llvm::Type *type) {
-  return impl->translateType(type);
-}

mlir/test/lib/Transforms/CMakeLists.txt

add_mlir_library(MLIRTestTransforms
   MLIRAffine
   MLIRAnalysis
   MLIREDSC
   MLIRGPU
   MLIRGPU
   MLIRGPUToGPURuntimeTransforms
   MLIRLinalg
   MLIRLinalgTransforms
+  MLIRLLVMToLLVMIRTranslation
   MLIRMathTransforms
   MLIRNVVMIR
   MLIRNVVMToLLVMIRTranslation
   MLIRPass
   MLIRROCDLIR
   MLIRROCDLToLLVMIRTranslation
   MLIRSCF
   MLIRSCFTransforms
   MLIRStandardOpsTransforms
-  MLIRTargetLLVMIR
+  MLIRTargetLLVMIRExport
   MLIRTestDialect
   MLIRTransformUtils
   MLIRVector
   MLIRVectorToSCF
   )
 
 include_directories(${CMAKE_CURRENT_SOURCE_DIR}/../Dialect/Test)
 include_directories(${CMAKE_CURRENT_BINARY_DIR}/../Dialect/Test)

mlir/test/lib/Transforms/TestConvertGPUKernelToCubin.cpp

 //===- TestConvertGPUKernelToCubin.cpp - Test gpu kernel cubin lowering ---===//
 //
 // 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
 //
 //===----------------------------------------------------------------------===//
 
 #include "mlir/Conversion/GPUCommon/GPUCommonPass.h"
 #include "mlir/Dialect/LLVMIR/NVVMDialect.h"
 #include "mlir/Pass/Pass.h"
 #include "mlir/Pass/PassManager.h"
-#include "mlir/Target/LLVMIR.h"
+#include "mlir/Target/LLVMIR/Dialect/LLVMIR/LLVMToLLVMIRTranslation.h"
 #include "mlir/Target/LLVMIR/Dialect/NVVM/NVVMToLLVMIRTranslation.h"
 #include "mlir/Target/LLVMIR/Export.h"
 #include "llvm/Support/TargetSelect.h"
 
 using namespace mlir;
 
 #if MLIR_CUDA_CONVERSIONS_ENABLED
 static OwnedBlob compilePtxToCubinForTesting(const std::string &, Location,
                                              StringRef) {
   const char data[] = "CUBIN";
   return std::make_unique<std::vector<char>>(data, data + sizeof(data) - 1);
 }
 
-static void registerNVVMDialectTranslation(MLIRContext &context) {
-  if (auto *dialect = context.getLoadedDialect<NVVM::NVVMDialect>()) {
-    if (!dialect->getRegisteredInterface<
-            NVVMDialectLLVMIRTranslationInterface>()) {
-      DialectRegistry registry;
-      registry.insert<NVVM::NVVMDialect>();
-      registry.addDialectInterface<NVVM::NVVMDialect,
-                                   NVVMDialectLLVMIRTranslationInterface>();
-      context.appendDialectRegistry(registry);
-    }
-  }
-}
-
 static std::unique_ptr<llvm::Module>
 translateModuleToNVVMIR(Operation *m, llvm::LLVMContext &llvmContext,
                         StringRef moduleName) {
   registerLLVMDialectTranslation(*m->getContext());
   registerNVVMDialectTranslation(*m->getContext());
   return translateModuleToLLVMIR(m, llvmContext, moduleName);
 }
 

mlir/test/lib/Transforms/TestConvertGPUKernelToHsaco.cpp

 //===- TestConvertGPUKernelToHsaco.cpp - Test gpu kernel hsaco lowering ---===//
 //
 // 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
 //
 //===----------------------------------------------------------------------===//
 
 #include "mlir/Conversion/GPUCommon/GPUCommonPass.h"
 #include "mlir/Dialect/LLVMIR/ROCDLDialect.h"
 #include "mlir/Pass/Pass.h"
 #include "mlir/Pass/PassManager.h"
-#include "mlir/Target/LLVMIR.h"
+#include "mlir/Target/LLVMIR/Dialect/LLVMIR/LLVMToLLVMIRTranslation.h"
 #include "mlir/Target/LLVMIR/Dialect/ROCDL/ROCDLToLLVMIRTranslation.h"
 #include "mlir/Target/LLVMIR/Export.h"
 #include "llvm/Support/TargetSelect.h"
 
 using namespace mlir;
 
 #if MLIR_ROCM_CONVERSIONS_ENABLED
 static OwnedBlob compileIsaToHsacoForTesting(const std::string &, Location,
                                              StringRef) {
   const char data[] = "HSACO";
   return std::make_unique<std::vector<char>>(data, data + sizeof(data) - 1);
 }
 
-static void registerROCDLDialectTranslation(MLIRContext &context) {
-  DialectRegistry registry;
-  registry.insert<ROCDL::ROCDLDialect>();
-  registry.addDialectInterface<ROCDL::ROCDLDialect,
-                               ROCDLDialectLLVMIRTranslationInterface>();
-  context.appendDialectRegistry(registry);
-}
-
 static std::unique_ptr<llvm::Module>
 translateModuleToROCDL(Operation *m, llvm::LLVMContext &llvmContext,
                        StringRef moduleName) {
   registerLLVMDialectTranslation(*m->getContext());
   registerROCDLDialectTranslation(*m->getContext());
   return translateModuleToLLVMIR(m, llvmContext, moduleName);
 }
 

mlir/tools/mlir-cpu-runner/CMakeLists.txt

 llvm_update_compile_flags(mlir-cpu-runner)
 target_link_libraries(mlir-cpu-runner PRIVATE
   MLIRAnalysis
   MLIREDSC
   MLIRExecutionEngine
   MLIRIR
   MLIRJitRunner
   MLIRLLVMIR
+  MLIRLLVMToLLVMIRTranslation
   MLIROpenMP
+  MLIROpenMPToLLVMIRTranslation
   MLIRParser
-  MLIRTargetLLVMIR
+  MLIRTargetLLVMIRExport
   MLIRSupport
   )

mlir/tools/mlir-cpu-runner/mlir-cpu-runner.cpp

 //===- mlir-cpu-runner.cpp - MLIR CPU Execution Driver---------------------===//
 //
 // 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
 //
 //===----------------------------------------------------------------------===//
 //
 // Main entry point to a command line utility that executes an MLIR file on the
 // CPU by  translating MLIR to LLVM IR before JIT-compiling and executing the
 // latter.
 //
 //===----------------------------------------------------------------------===//
 
 #include "mlir/Dialect/LLVMIR/LLVMDialect.h"
-#include "mlir/Dialect/OpenMP/OpenMPDialect.h"
 #include "mlir/ExecutionEngine/JitRunner.h"
 #include "mlir/ExecutionEngine/OptUtils.h"
 #include "mlir/IR/Dialect.h"
-#include "mlir/Target/LLVMIR.h"
+#include "mlir/Target/LLVMIR/Dialect/LLVMIR/LLVMToLLVMIRTranslation.h"
 #include "mlir/Target/LLVMIR/Dialect/OpenMP/OpenMPToLLVMIRTranslation.h"
 
 #include "llvm/Support/InitLLVM.h"
 #include "llvm/Support/TargetSelect.h"
 
 int main(int argc, char **argv) {
   llvm::InitLLVM y(argc, argv);
   llvm::InitializeNativeTarget();
   llvm::InitializeNativeTargetAsmPrinter();
   llvm::InitializeNativeTargetAsmParser();
   mlir::initializeLLVMPasses();
 
   mlir::DialectRegistry registry;
-  registry.insert<mlir::LLVM::LLVMDialect, mlir::omp::OpenMPDialect>();
   mlir::registerLLVMDialectTranslation(registry);
-  registry.addDialectInterface<mlir::omp::OpenMPDialect,
-                               mlir::OpenMPDialectLLVMIRTranslationInterface>();
+  mlir::registerOpenMPDialectTranslation(registry);
 
   return mlir::JitRunnerMain(argc, argv, registry);
 }

mlir/tools/mlir-cuda-runner/CMakeLists.txt

set(LIBS
     MLIRJitRunner
     MLIRAnalysis
     MLIRAsync
     MLIREDSC
     MLIRExecutionEngine
     MLIRGPU
     MLIRIR
     MLIRLLVMIR
+    MLIRLLVMToLLVMIRTranslation
     MLIRNVVMIR
     MLIRParser
     MLIRStandard
     MLIRSupport
-    MLIRTargetLLVMIR
+    MLIRTargetLLVMIRExport
     MLIRNVVMToLLVMIRTranslation
     MLIRTransforms
     MLIRTranslation
     ${CUDA_RUNTIME_LIBRARY}
   )
 
   # Manually expand the target library, since our MLIR libraries
   # aren't plugged into the LLVM dependency tracking. If we don't

mlir/tools/mlir-cuda-runner/mlir-cuda-runner.cpp

 #include "mlir/Dialect/LLVMIR/LLVMDialect.h"
 #include "mlir/Dialect/LLVMIR/NVVMDialect.h"
 #include "mlir/Dialect/StandardOps/IR/Ops.h"
 #include "mlir/ExecutionEngine/JitRunner.h"
 #include "mlir/ExecutionEngine/OptUtils.h"
 #include "mlir/IR/BuiltinOps.h"
 #include "mlir/Pass/Pass.h"
 #include "mlir/Pass/PassManager.h"
-#include "mlir/Target/LLVMIR.h"
+#include "mlir/Target/LLVMIR/Dialect/LLVMIR/LLVMToLLVMIRTranslation.h"
 #include "mlir/Target/LLVMIR/Dialect/NVVM/NVVMToLLVMIRTranslation.h"
 #include "mlir/Target/LLVMIR/Export.h"
 #include "mlir/Transforms/DialectConversion.h"
 #include "mlir/Transforms/Passes.h"
 
 #include "llvm/Support/InitLLVM.h"
 #include "llvm/Support/TargetSelect.h"
 
int main(int argc, char **argv) {
 
   mlir::JitRunnerConfig jitRunnerConfig;
   jitRunnerConfig.mlirTransformer = mlirTransformer;
 
   mlir::DialectRegistry registry;
   registry.insert<mlir::LLVM::LLVMDialect, mlir::NVVM::NVVMDialect,
                   mlir::async::AsyncDialect, mlir::gpu::GPUDialect,
                   mlir::StandardOpsDialect>();
-  registry.addDialectInterface<NVVM::NVVMDialect,
-                               NVVMDialectLLVMIRTranslationInterface>();
   mlir::registerLLVMDialectTranslation(registry);
+  mlir::registerNVVMDialectTranslation(registry);
 
   return mlir::JitRunnerMain(argc, argv, registry, jitRunnerConfig);
 }

mlir/tools/mlir-rocm-runner/CMakeLists.txt

set(LIBS
     lldELF
     MLIRJitRunner
     MLIRAnalysis
     MLIREDSC
     MLIRExecutionEngine
     MLIRGPU
     MLIRIR
     MLIRLLVMIR
+    MLIRLLVMToLLVMIRTranslation
     MLIRParser
     MLIRROCDLIR
     MLIRStandard
     MLIRSupport
-    MLIRTargetLLVMIR
+    MLIRTargetLLVMIRExport
     MLIRROCDLToLLVMIRTranslation
     MLIRTransforms
     MLIRTranslation
     ${ROCM_RUNTIME_LIBRARY}
   )
 
   # Manually expand the target library, since our MLIR libraries
   # aren't plugged into the LLVM dependency tracking. If we don't

mlir/tools/mlir-rocm-runner/mlir-rocm-runner.cpp

 #include "mlir/Dialect/LLVMIR/ROCDLDialect.h"
 #include "mlir/Dialect/StandardOps/IR/Ops.h"
 #include "mlir/ExecutionEngine/JitRunner.h"
 #include "mlir/ExecutionEngine/OptUtils.h"
 #include "mlir/IR/BuiltinOps.h"
 #include "mlir/Pass/Pass.h"
 #include "mlir/Pass/PassManager.h"
 #include "mlir/Support/FileUtilities.h"
-#include "mlir/Target/LLVMIR.h"
+#include "mlir/Target/LLVMIR/Dialect/LLVMIR/LLVMToLLVMIRTranslation.h"
 #include "mlir/Target/LLVMIR/Dialect/ROCDL/ROCDLToLLVMIRTranslation.h"
 #include "mlir/Target/LLVMIR/Export.h"
 #include "mlir/Transforms/DialectConversion.h"
 #include "mlir/Transforms/Passes.h"
 #include "llvm/Support/ErrorOr.h"
 #include "llvm/Support/FileUtilities.h"
 #include "llvm/Support/InitLLVM.h"
 #include "llvm/Support/LineIterator.h"
int main(int argc, char **argv) {
   mlir::initializeLLVMPasses();
 
   mlir::JitRunnerConfig jitRunnerConfig;
   jitRunnerConfig.mlirTransformer = runMLIRPasses;
 
   mlir::DialectRegistry registry;
   registry.insert<mlir::LLVM::LLVMDialect, mlir::gpu::GPUDialect,
                   mlir::ROCDL::ROCDLDialect, mlir::StandardOpsDialect>();
-  registry.addDialectInterface<ROCDL::ROCDLDialect,
-                               ROCDLDialectLLVMIRTranslationInterface>();
   mlir::registerLLVMDialectTranslation(registry);
+  mlir::registerROCDLDialectTranslation(registry);
 
   return mlir::JitRunnerMain(argc, argv, registry, jitRunnerConfig);
 }

mlir/tools/mlir-spirv-cpu-runner/CMakeLists.txt

target_link_libraries(mlir-spirv-cpu-runner PRIVATE
     ${conversion_libs}
     MLIRAnalysis
     MLIREDSC
     MLIRExecutionEngine
     MLIRGPU
     MLIRIR
     MLIRJitRunner
     MLIRLLVMIR
+    MLIRLLVMToLLVMIRTranslation
     MLIRParser
     MLIRSPIRV
     MLIRStandard
-    MLIRTargetLLVMIR
+    MLIRTargetLLVMIRExport
     MLIRTransforms
     MLIRTranslation
     MLIRSupport
   )
 endif()

mlir/tools/mlir-spirv-cpu-runner/mlir-spirv-cpu-runner.cpp

 #include "mlir/Dialect/SPIRV/IR/SPIRVDialect.h"
 #include "mlir/Dialect/SPIRV/IR/SPIRVOps.h"
 #include "mlir/Dialect/SPIRV/Transforms/Passes.h"
 #include "mlir/Dialect/StandardOps/IR/Ops.h"
 #include "mlir/ExecutionEngine/JitRunner.h"
 #include "mlir/ExecutionEngine/OptUtils.h"
 #include "mlir/Pass/Pass.h"
 #include "mlir/Pass/PassManager.h"
-#include "mlir/Target/LLVMIR.h"
+#include "mlir/Target/LLVMIR/Dialect/LLVMIR/LLVMToLLVMIRTranslation.h"
 #include "mlir/Target/LLVMIR/Export.h"
 
 #include "llvm/IR/LLVMContext.h"
 #include "llvm/IR/Module.h"
 #include "llvm/Linker/Linker.h"
 #include "llvm/Support/InitLLVM.h"
 #include "llvm/Support/TargetSelect.h"
 

mlir/tools/mlir-vulkan-runner/CMakeLists.txt

set(LIBS
     ${conversion_libs}
     MLIRAnalysis
     MLIREDSC
     MLIRExecutionEngine
     MLIRGPU
     MLIRIR
     MLIRJitRunner
     MLIRLLVMIR
+    MLIRLLVMToLLVMIRTranslation
     MLIRParser
     MLIRSPIRV
     MLIRSPIRVTransforms
     MLIRStandard
     MLIRSupport
-    MLIRTargetLLVMIR
+    MLIRTargetLLVMIRExport
     MLIRTransforms
     MLIRTranslation
     ${Vulkan_LIBRARY}
   )
 
   # Manually expand the target library, since our MLIR libraries
   # aren't plugged into the LLVM dependency tracking. If we don't
   # do this then we can't insert the CodeGen library after ourselves

mlir/tools/mlir-vulkan-runner/mlir-vulkan-runner.cpp

 #include "mlir/Dialect/SPIRV/IR/SPIRVDialect.h"
 #include "mlir/Dialect/SPIRV/IR/SPIRVOps.h"
 #include "mlir/Dialect/SPIRV/Transforms/Passes.h"
 #include "mlir/Dialect/StandardOps/IR/Ops.h"
 #include "mlir/ExecutionEngine/JitRunner.h"
 #include "mlir/ExecutionEngine/OptUtils.h"
 #include "mlir/Pass/Pass.h"
 #include "mlir/Pass/PassManager.h"
-#include "mlir/Target/LLVMIR.h"
+#include "mlir/Target/LLVMIR/Dialect/LLVMIR/LLVMToLLVMIRTranslation.h"
+#include "mlir/Target/LLVMIR/Export.h"
 #include "llvm/Support/InitLLVM.h"
 #include "llvm/Support/TargetSelect.h"
 
 using namespace mlir;
 
 static LogicalResult runMLIRPasses(ModuleOp module) {
   PassManager passManager(module.getContext());
   applyPassManagerCLOptions(passManager);

mlir/unittests/ExecutionEngine/Invoke.cpp

 #include "mlir/ExecutionEngine/CRunnerUtils.h"
 #include "mlir/ExecutionEngine/ExecutionEngine.h"
 #include "mlir/ExecutionEngine/MemRefUtils.h"
 #include "mlir/ExecutionEngine/RunnerUtils.h"
 #include "mlir/IR/MLIRContext.h"
 #include "mlir/InitAllDialects.h"
 #include "mlir/Parser.h"
 #include "mlir/Pass/PassManager.h"
-#include "mlir/Target/LLVMIR.h"
+#include "mlir/Target/LLVMIR/Dialect/LLVMIR/LLVMToLLVMIRTranslation.h"
+#include "mlir/Target/LLVMIR/Export.h"
 #include "llvm/Support/TargetSelect.h"
 #include "llvm/Support/raw_ostream.h"
 
 #include "gmock/gmock.h"
 
 using namespace mlir;
 
 static struct LLVMInitializer {
TEST(NativeMemRefJit, JITCallback) {
   int64_t shape[] = {K, M};
   int64_t shapeAlloc[] = {K + 1, M + 1};
   OwningMemRef<float, 2> A(shape, shapeAlloc);
   int count = 1;
   for (float &elt : *A)
     elt = count++;
 
   std::string moduleStr = R"mlir(
-  func private @callback(%arg0: memref<?x?xf32>, %coefficient: i32)  attributes { llvm.emit_c_interface } 
+  func private @callback(%arg0: memref<?x?xf32>, %coefficient: i32)  attributes { llvm.emit_c_interface }
   func @caller_for_callback(%arg0: memref<?x?xf32>, %coefficient: i32) attributes { llvm.emit_c_interface } {
     %unranked = memref_cast %arg0: memref<?x?xf32> to memref<*xf32>
     call @callback(%arg0, %coefficient) : (memref<?x?xf32>, i32) -> ()
     return
   }
   )mlir";
   DialectRegistry registry;
   registerAllDialects(registry);