[mlir] Use the interface-based translation for LLVM "intrinsic" dialects

[ftynse]

Summary

Port the translation of five dialects that define LLVM IR intrinsics
(LLVMAVX512, LLVMArmNeon, LLVMArmSVE, NVVM, ROCDL) to the new dialect
interface-based mechanism. This allows us to remove individual translations
that were created for each of these dialects and just use one common
MLIR-to-LLVM-IR translation that potentially supports all dialects instead,
based on what is registered and including any combination of translatable
dialects. This removal was one of the main goals of the refactoring.

To support the addition of GPU-related metadata, the translation interface is
extended with the amendOperation function that allows the interface
implementation to post-process any translated operation with dialect attributes
from the dialect for which the interface is implemented regardless of the
operation's dialect. This is currently applied to "kernel" functions, but can
be used to construct other metadata in dialect-specific ways without
necessarily affecting operations.

Depends On D96591, D96504

[aartbik]

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

21

typo: belonging

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

21

ditto

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

21

and one more, and below, won't repeat

[ftynse]

Rebase and configure build

[nicolasvasilache]

mlir/lib/Target/LLVMIR/ConvertToLLVMIR.cpp

83–85

Do we intend to pull the kitchen sink in the longer term or is it just transient state waiting for e.g. better opt-in CMake support?

[ftynse]

mlir/lib/Target/LLVMIR/ConvertToLLVMIR.cpp

83–85

I'll refactor it a bit further to have a mlir::registerAllLLVMIRTranslations by analogy with mlir::registerAllDialects used in mlir-opt, but we will have to have explicit registration somewhere unless we want link-time registration, which would contradict the rest of MLIR.

[mehdi_amini]

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

53

Should this be in a verifier?

mlir/test/Target/avx512.mlir

1

Was this totally unnecessary ? We have coverage for the round-trip of whatever is here?

[ftynse]

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

53

9cd47a26d593985ad2b36857cb9fcbc7659ffde8

mlir/test/Target/avx512.mlir

1

There is no custom syntax for these ops, and there are no diagnostics checks either. Otherwise, I would have split out the roundtripping part anyway, "target" tests should not test it.

mlir/examples/standalone/test/Standalone/standalone-translate.mlir

 // RUN: standalone-translate --help | FileCheck %s
-// CHECK: --avx512-mlir-to-llvmir
 // CHECK: --deserialize-spirv
 // CHECK: --import-llvm
 // CHECK: --mlir-to-llvmir
-// CHECK: --mlir-to-nvvmir
-// CHECK: --mlir-to-rocdlir
 // CHECK: --serialize-spirv

mlir/include/mlir/Dialect/LLVMIR/LLVMOpBase.td

list<string> lst = !foreach(x, values,
                               !subst("$0", !cast<string>(x), pattern));
 }
 
 // Patterns with code obtaining the LLVM IR type of the given operand or result
 // of operation. "$0" is expected to be replaced by the position of the operand
 // or result in the operation.
 def LLVM_IntrPatterns {
   string operand =
-    [{convertType(opInst.getOperand($0).getType())}];
+    [{moduleTranslation.convertType(opInst.getOperand($0).getType())}];
   string result =
-    [{convertType(opInst.getResult($0).getType())}];
+    [{moduleTranslation.convertType(opInst.getResult($0).getType())}];
   string structResult =
-    [{convertType(opInst.getResult(0).getType().cast<LLVM::LLVMStructType>()
-                                                   .getBody()[$0])}];
+    [{moduleTranslation.convertType(
+        opInst.getResult(0).getType().cast<LLVM::LLVMStructType>()
 }
 
 
 // Base class for LLVM intrinsics operation. It is similar to LLVM_Op, but
 // provides the "llvmBuilder" field for constructing the intrinsic.
 // The builder relies on the contents of "overloadedResults" and
 // "overloadedOperands" lists that contain the positions of intrinsic results
 // and operands that are overloadable in the LLVM sense, that is their types
let llvmBuilder = [{
     llvm::Function *fn = llvm::Intrinsic::getDeclaration(
         module,
         llvm::Intrinsic::}] # enumName # [{,
         { }] # !interleave(!listconcat(
             ListIntSubst<resultPattern, overloadedResults>.lst,
             ListIntSubst<LLVM_IntrPatterns.operand,
                          overloadedOperands>.lst), ", ") # [{
         });
-    auto operands = lookupValues(opInst.getOperands());
+    auto operands = moduleTranslation.lookupValues(opInst.getOperands());
     }] # !if(!gt(numResults, 0), "$res = ", "")
        # [{builder.CreateCall(fn, operands);
   }];
 }
 
 // Base class for LLVM intrinsic operations, should not be used directly. Places
 // the intrinsic into the LLVM dialect and prefixes its name with "intr.".
 class LLVM_IntrOp<string mnem, list<int> overloadedResults,
class LLVM_VectorReductionAcc<string mnem>
   let llvmBuilder = [{
     llvm::Module *module = builder.GetInsertBlock()->getModule();
     llvm::Function *fn = llvm::Intrinsic::getDeclaration(
         module,
         llvm::Intrinsic::vector_reduce_}] # mnem # [{,
         { }] # !interleave(ListIntSubst<LLVM_IntrPatterns.operand, [1]>.lst,
                            ", ") # [{
         });
-    auto operands = lookupValues(opInst.getOperands());
+    auto operands = moduleTranslation.lookupValues(opInst.getOperands());
     llvm::FastMathFlags origFM = builder.getFastMathFlags();
     llvm::FastMathFlags tempFM = origFM;
     tempFM.setAllowReassoc($reassoc);
     builder.setFastMathFlags(tempFM);  // set fastmath flag
     $res = builder.CreateCall(fn, operands);
     builder.setFastMathFlags(origFM);  // restore fastmath flag
   }];
 }
 
 #endif  // LLVMIR_OP_BASE

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

def LLVM_UndefOp : LLVM_Op<"mlir.undef", [NoSideEffect]>,
   }];
   let results = (outs LLVM_Type:$res);
   let builders = [LLVM_OneResultOpBuilder];
   let assemblyFormat = "attr-dict `:` type($res)";
 }
 
 def LLVM_ConstantOp
     : LLVM_Op<"mlir.constant", [NoSideEffect]>,
-      LLVM_Builder<"$res = getLLVMConstant($_resultType, $value, $_location);">
+      LLVM_Builder<[{$res = getLLVMConstant($_resultType, $value, $_location,
+                                            moduleTranslation);}]>
 {
   let summary = "Defines a constant of LLVM type.";
   let description = [{
     Unlike LLVM IR, MLIR does not have first-class constant values. Therefore,
     all constants must be created as SSA values before being used in other
     operations. `llvm.mlir.constant` creates such values for scalars and
     vectors. It has a mandatory `value` attribute, which may be an integer,
     floating point attribute; dense or sparse attribute containing integers or

mlir/include/mlir/Dialect/LLVMIR/ROCDLOps.td

def ROCDL_MubufStoreOp :
   ROCDL_Op<"buffer.store">,
   Arguments<(ins LLVM_Type:$vdata,
                  LLVM_Type:$rsrc,
                  LLVM_Type:$vindex,
                  LLVM_Type:$offset,
                  LLVM_Type:$glc,
                  LLVM_Type:$slc)>{
   string llvmBuilder = [{
-    auto vdataType = convertType(op.vdata().getType());
+    auto vdataType = moduleTranslation.convertType(op.vdata().getType());
     createIntrinsicCall(builder,
           llvm::Intrinsic::amdgcn_buffer_store, {$vdata, $rsrc, $vindex,
           $offset, $glc, $slc}, {vdataType});
   }];
   let parser = [{ return parseROCDLMubufStoreOp(parser, result); }];
   let printer = [{
     Operation *op = this->getOperation();
     p << op->getName() << " " << op->getOperands()
       << " : " << vdata().getType();
   }];
 }
 
 #endif // ROCDLIR_OPS

mlir/include/mlir/InitAllTranslations.h

 #define MLIR_INITALLTRANSLATIONS_H
 
 namespace mlir {
 
 void registerFromLLVMIRTranslation();
 void registerFromSPIRVTranslation();
 void registerToLLVMIRTranslation();
 void registerToSPIRVTranslation();
-void registerToNVVMIRTranslation();
-void registerToROCDLIRTranslation();
-void registerArmNeonToLLVMIRTranslation();
-void registerAVX512ToLLVMIRTranslation();
-void registerArmSVEToLLVMIRTranslation();
 
 // This function should be called before creating any MLIRContext if one
 // expects all the possible translations to be made available to the context
 // automatically.
 inline void registerAllTranslations() {
   static bool initOnce = []() {
     registerFromLLVMIRTranslation();
     registerFromSPIRVTranslation();
     registerToLLVMIRTranslation();
     registerToSPIRVTranslation();
-    registerToNVVMIRTranslation();
-    registerToROCDLIRTranslation();
-    registerArmNeonToLLVMIRTranslation();
-    registerAVX512ToLLVMIRTranslation();
-    registerArmSVEToLLVMIRTranslation();
     return true;
   }();
   (void)initOnce;
 }
 } // namespace mlir
 
 #endif // MLIR_INITALLTRANSLATIONS_H

mlir/include/mlir/Target/LLVMIR.h

 class Module;
 } // namespace llvm
 
 namespace mlir {
 
 class DialectRegistry;
 class OwningModuleRef;
 class MLIRContext;
-class ModuleOp;
+class Operation;
 
 /// Convert the given MLIR module into LLVM IR.  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 `nullptr`.
 std::unique_ptr<llvm::Module>
-translateModuleToLLVMIR(ModuleOp m, llvm::LLVMContext &llvmContext,
+translateModuleToLLVMIR(Operation *op, llvm::LLVMContext &llvmContext,
                         StringRef name = "LLVMDialectModule");
 
 /// 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,

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

-//===- LLVMToLLVMIRTranslation.h - LLVM Dialect to LLVM IR-------*- C++ -*-===//
+//===- 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 LLVM dialect
-// to LLVM IR.
+// This file implements the dialect interface for translating the LLVMAVX512
+// dialect to LLVM IR.
 //
 //===----------------------------------------------------------------------===//
 
-#ifndef MLIR_TARGET_LLVMIR_DIALECT_LLVMIR_LLVMTOLLVMIRTRANSLATION_H
-#define MLIR_TARGET_LLVMIR_DIALECT_LLVMIR_LLVMTOLLVMIRTRANSLATION_H
+#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 beloning to
-/// the LLVM dialect to LLVM IR.
-class LLVMDialectLLVMIRTranslationInterface
+/// 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;
 };
 
 } // namespace mlir
 
-#endif // MLIR_TARGET_LLVMIR_DIALECT_LLVMIR_LLVMTOLLVMIRTRANSLATION_H
+#endif // MLIR_TARGET_LLVMIR_DIALECT_LLVMAVX512_LLVMAVX512TOLLVMIRTRANSLATION_H

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

-//===- LLVMToLLVMIRTranslation.h - LLVM Dialect to LLVM IR-------*- C++ -*-===//
+//===- 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 LLVM dialect
-// to LLVM IR.
+// This file implements the dialect interface for translating the LLVMArmNeon
+// dialect to LLVM IR.
 //
 //===----------------------------------------------------------------------===//
 
-#ifndef MLIR_TARGET_LLVMIR_DIALECT_LLVMIR_LLVMTOLLVMIRTRANSLATION_H
-#define MLIR_TARGET_LLVMIR_DIALECT_LLVMIR_LLVMTOLLVMIRTRANSLATION_H
+#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 beloning to
-/// the LLVM dialect to LLVM IR.
-class LLVMDialectLLVMIRTranslationInterface
+/// 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;
 };
 
 } // namespace mlir
 
-#endif // MLIR_TARGET_LLVMIR_DIALECT_LLVMIR_LLVMTOLLVMIRTRANSLATION_H
+#endif // MLIR_TARGET_LLVMIR_DIALECT_LLVMARMNEON_LLVMARMNEONTOLLVMIRTRANSLATION_H

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

-//===- LLVMToLLVMIRTranslation.h - LLVM Dialect to LLVM IR-------*- C++ -*-===//
+//===- 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 LLVM dialect
-// to LLVM IR.
+// This file implements the dialect interface for translating the LLVMArmSVE
+// dialect to LLVM IR.
 //
 //===----------------------------------------------------------------------===//
 
-#ifndef MLIR_TARGET_LLVMIR_DIALECT_LLVMIR_LLVMTOLLVMIRTRANSLATION_H
-#define MLIR_TARGET_LLVMIR_DIALECT_LLVMIR_LLVMTOLLVMIRTRANSLATION_H
+#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 beloning to
-/// the LLVM dialect to LLVM IR.
-class LLVMDialectLLVMIRTranslationInterface
+/// 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;
 };
 
 } // namespace mlir
 
-#endif // MLIR_TARGET_LLVMIR_DIALECT_LLVMIR_LLVMTOLLVMIRTRANSLATION_H
+#endif // MLIR_TARGET_LLVMIR_DIALECT_LLVMARMSVE_LLVMARMSVETOLLVMIRTRANSLATION_H

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

 
 #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 beloning to
-/// the LLVM dialect to LLVM IR.
+/// 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;
 };
 
 } // namespace mlir
 
 #endif // MLIR_TARGET_LLVMIR_DIALECT_LLVMIR_LLVMTOLLVMIRTRANSLATION_H

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

-//===- LLVMToLLVMIRTranslation.h - LLVM Dialect to LLVM IR-------*- C++ -*-===//
+//===- 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 LLVM dialect
-// to LLVM IR.
+// This file implements the dialect interface for translating the NVVM
+// dialect to LLVM IR.
 //
 //===----------------------------------------------------------------------===//
 
-#ifndef MLIR_TARGET_LLVMIR_DIALECT_LLVMIR_LLVMTOLLVMIRTRANSLATION_H
-#define MLIR_TARGET_LLVMIR_DIALECT_LLVMIR_LLVMTOLLVMIRTRANSLATION_H
+#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 beloning to
-/// the LLVM dialect to LLVM IR.
-class LLVMDialectLLVMIRTranslationInterface
+/// 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;
 };
 
 } // namespace mlir
 
-#endif // MLIR_TARGET_LLVMIR_DIALECT_LLVMIR_LLVMTOLLVMIRTRANSLATION_H
+#endif // MLIR_TARGET_LLVMIR_DIALECT_NVVM_NVVMTOLLVMIRTRANSLATION_H

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

 
 #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 beloning to
-/// the OpenMP dialect to LLVM IR.
+/// 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;
 };
 
 } // namespace mlir
 
 #endif // MLIR_TARGET_LLVMIR_DIALECT_OPENMP_OPENMPTOLLVMIRTRANSLATION_H

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

-//===- LLVMToLLVMIRTranslation.h - LLVM Dialect to LLVM IR-------*- C++ -*-===//
+//===- 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 LLVM dialect
-// to LLVM IR.
+// This file implements the dialect interface for translating the ROCDL
+// dialect to LLVM IR.
 //
 //===----------------------------------------------------------------------===//
 
-#ifndef MLIR_TARGET_LLVMIR_DIALECT_LLVMIR_LLVMTOLLVMIRTRANSLATION_H
-#define MLIR_TARGET_LLVMIR_DIALECT_LLVMIR_LLVMTOLLVMIRTRANSLATION_H
+#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 beloning to
-/// the LLVM dialect to LLVM IR.
-class LLVMDialectLLVMIRTranslationInterface
+/// 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;
 };
 
 } // namespace mlir
 
-#endif // MLIR_TARGET_LLVMIR_DIALECT_LLVMIR_LLVMTOLLVMIRTRANSLATION_H
+#endif // MLIR_TARGET_LLVMIR_DIALECT_ROCDL_ROCDLTOLLVMIRTRANSLATION_H

mlir/include/mlir/Target/LLVMIR/LLVMTranslationInterface.h

 //===- LLVMTranslationInterface.h - Translation to LLVM iface ---*- 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 header file defines dialect interfaces for translation to LLVM IR.
 //
 //===----------------------------------------------------------------------===//
 
 #ifndef MLIR_TARGET_LLVMIR_LLVMTRANSLATIONINTERFACE_H
 #define MLIR_TARGET_LLVMIR_LLVMTRANSLATIONINTERFACE_H
 
+#include "mlir/IR/Attributes.h"
 #include "mlir/IR/DialectInterface.h"
+#include "mlir/IR/Identifier.h"
 #include "mlir/Support/LogicalResult.h"
 
 namespace llvm {
 class IRBuilderBase;
-}
+} // namespace llvm
 
 namespace mlir {
 namespace LLVM {
 class ModuleTranslation;
 } // namespace LLVM
 
 /// Base class for dialect interfaces providing translation to LLVM IR.
 /// Dialects that can be translated should provide an implementation of this
 /// interface for the supported operations. The interface may be implemented in
 /// a separate library to avoid the "main" dialect library depending on LLVM IR.
 /// The interface can be attached using the delayed registration mechanism
 /// available in DialectRegistry.
 class LLVMTranslationDialectInterface
     : public DialectInterface::Base<LLVMTranslationDialectInterface> {
 public:
   LLVMTranslationDialectInterface(Dialect *dialect) : Base(dialect) {}
 
   /// Hook for derived dialect interface to provide translation of the
   /// operations to LLVM IR.
   virtual LogicalResult
   convertOperation(Operation *op, llvm::IRBuilderBase &builder,
                    LLVM::ModuleTranslation &moduleTranslation) const {
     return failure();
   }
+
+  /// Hook for derived dialect interface to act on an operation that has dialect
+  /// attributes from the derived dialect (the operation itself may be from a
+  /// different dialect). This gets called after the operation has been
+  /// translated. The hook is expected to use moduleTranslation to look up the
+  /// translation results and amend the corresponding IR constructs. Does
+  /// nothing and succeeds by default.
+  virtual LogicalResult
+  amendOperation(Operation *op, NamedAttribute attribute,
+                 LLVM::ModuleTranslation &moduleTranslation) const {
+    return success();
+  }
 };
 
 /// Interface collection for translation to LLVM IR, dispatches to a concrete
 /// interface implementation based on the dialect to which the given op belongs.
 class LLVMTranslationInterface
     : public DialectInterfaceCollection<LLVMTranslationDialectInterface> {
 public:
   using Base::Base;
 
   /// Translates the given operation to LLVM IR using the interface implemented
   /// by the op's dialect.
   virtual LogicalResult
   convertOperation(Operation *op, llvm::IRBuilderBase &builder,
                    LLVM::ModuleTranslation &moduleTranslation) const {
     if (const LLVMTranslationDialectInterface *iface = getInterfaceFor(op))
       return iface->convertOperation(op, builder, moduleTranslation);
     return failure();
   }
+
+  /// Acts on the given operation using the interface implemented by the dialect
+  /// of one of the operation's dialect attributes.
+  virtual LogicalResult
+  amendOperation(Operation *op, NamedAttribute attribute,
+                 LLVM::ModuleTranslation &moduleTranslation) const {
+    if (const LLVMTranslationDialectInterface *iface =
+            getInterfaceFor(attribute.first.getDialect())) {
+      return iface->amendOperation(op, attribute, moduleTranslation);
+    }
+    return success();
+  }
 };
 
 } // namespace mlir
 
 #endif // MLIR_TARGET_LLVMIR_LLVMTRANSLATIONINTERFACE_H

mlir/include/mlir/Target/LLVMIR/ModuleTranslation.h

public:
   }
 
   /// Converts the type from MLIR LLVM dialect to LLVM.
   llvm::Type *convertType(Type type);
 
   /// Looks up remapped a list of remapped values.
   SmallVector<llvm::Value *, 8> lookupValues(ValueRange values);
 
-  /// Create an LLVM IR constant of `llvmType` from the MLIR attribute `attr`.
-  /// This currently supports integer, floating point, splat and dense element
-  /// attributes and combinations thereof.  In case of error, report it to `loc`
-  /// and return nullptr.
-  llvm::Constant *getLLVMConstant(llvm::Type *llvmType, Attribute attr,
-                                  Location loc);
-
   /// Returns the MLIR context of the module being translated.
   MLIRContext &getContext() { return *mlirModule->getContext(); }
 
   /// Returns the LLVM context in which the IR is being constructed.
-  llvm::LLVMContext &getLLVMContext() { return llvmModule->getContext(); }
+  llvm::LLVMContext &getLLVMContext() const { return llvmModule->getContext(); }
 
   /// Finds an LLVM IR global value that corresponds to the given MLIR operation
   /// defining a global value.
   llvm::GlobalValue *lookupGlobal(Operation *op) {
     return globalsMapping.lookup(op);
   }
 
   /// Returns the OpenMP IR builder associated with the LLVM IR module being
public:
   /// expected to exist in the mapping of this translator. Uses `builder` to
   /// translate the IR, leaving it at the end of the block. If `ignoreArguments`
   /// is set, does not produce PHI nodes for the block arguments. Otherwise, the
   /// PHI nodes are constructed for block arguments but are _not_ connected to
   /// the predecessors that may not exist yet.
   LogicalResult convertBlock(Block &bb, bool ignoreArguments,
                              llvm::IRBuilder<> &builder);
 
+  /// Gets the named metadata in the LLVM IR module being constructed, creating
+  /// it if it does not exist.
+  llvm::NamedMDNode *getOrInsertNamedModuleMetadata(StringRef name);
+
 protected:
   /// Translate the given MLIR module expressed in MLIR LLVM IR dialect into an
   /// LLVM IR module. The MLIR LLVM IR dialect holds a pointer to an
   /// LLVMContext, the LLVM IR module will be created in that context.
   ModuleTranslation(Operation *module,
                     std::unique_ptr<llvm::Module> llvmModule);
   virtual ~ModuleTranslation();
 
   virtual LogicalResult convertOperation(Operation &op,
                                          llvm::IRBuilder<> &builder);
 
   static std::unique_ptr<llvm::Module>
   prepareLLVMModule(Operation *m, llvm::LLVMContext &llvmContext,
                     StringRef name);
 
 private:
   /// Check whether the module contains only supported ops directly in its body.
   static LogicalResult checkSupportedModuleOps(Operation *m);
 
   LogicalResult convertFunctionSignatures();
   LogicalResult convertFunctions();
   LogicalResult convertGlobals();
   LogicalResult convertOneFunction(LLVMFuncOp func);
 
+  /// Translates dialect attributes attached to the given operation.
+  LogicalResult convertDialectAttributes(Operation *op);
+
   /// Original and translated module.
   Operation *mlirModule;
   std::unique_ptr<llvm::Module> llvmModule;
   /// A converter for translating debug information.
   std::unique_ptr<detail::DebugTranslation> debugTranslation;
 
   /// Builder for LLVM IR generation of OpenMP constructs.
   std::unique_ptr<llvm::OpenMPIRBuilder> ompBuilder;
 
   /// Precomputed pointer to OpenMP dialect. Note this can be nullptr if the
   /// OpenMP dialect hasn't been loaded (it is always loaded if there are OpenMP
   /// operations in the module though).
   const Dialect *ompDialect;
 
   /// Mappings between llvm.mlir.global definitions and corresponding globals.
   DenseMap<Operation *, llvm::GlobalValue *> globalsMapping;
 
   /// A stateful object used to translate types.
   TypeToLLVMIRTranslator typeTranslator;
 
+  /// A dialect interface collection used for dispatching the translation to
+  /// specific dialects.
   LLVMTranslationInterface iface;
 
   /// Mappings between original and translated values, used for lookups.
   llvm::StringMap<llvm::Function *> functionMapping;
   DenseMap<Value, llvm::Value *> valueMapping;
   DenseMap<Block *, llvm::BasicBlock *> blockMapping;
 
   /// A mapping between MLIR LLVM dialect terminators and LLVM IR terminators
   /// they are converted to. This allows for connecting PHI nodes to the source
   /// values after all operations are converted.
   DenseMap<Operation *, llvm::Instruction *> branchMapping;
 };
 
 namespace detail {
 /// For all blocks in the region that were converted to LLVM IR using the given
 /// ModuleTranslation, connect the PHI nodes of the corresponding LLVM IR blocks
 /// to the results of preceding blocks.
 void connectPHINodes(Region &region, const ModuleTranslation &state);
 
 /// Get a topologically sorted list of blocks of the given region.
 llvm::SetVector<Block *> getTopologicallySortedBlocks(Region &region);
+
+/// Create an LLVM IR constant of `llvmType` from the MLIR attribute `attr`.
+/// This currently supports integer, floating point, splat and dense element
+/// attributes and combinations thereof.  In case of error, report it to `loc`
+/// and return nullptr.
+llvm::Constant *getLLVMConstant(llvm::Type *llvmType, Attribute attr,
+                                Location loc,
+                                const ModuleTranslation &moduleTranslation);
+
+/// Creates a call to an LLVM IR intrinsic function with the given arguments.
+llvm::Value *createIntrinsicCall(llvm::IRBuilderBase &builder,
+                                 llvm::Intrinsic::ID intrinsic,
+                                 ArrayRef<llvm::Value *> args = {},
+                                 ArrayRef<llvm::Type *> tys = {});
 } // namespace detail
 
 } // namespace LLVM
 } // namespace mlir
 
 #endif // MLIR_TARGET_LLVMIR_MODULETRANSLATION_H

mlir/include/mlir/Target/NVVMIR.h

-//===- NVVMIR.h - MLIR to LLVM + NVVM IR conversion -------------*- 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 declares the entry point for the MLIR to LLVM + NVVM IR conversion.
-//
-//===----------------------------------------------------------------------===//
-
-#ifndef MLIR_TARGET_NVVMIR_H
-#define MLIR_TARGET_NVVMIR_H
-
-#include "llvm/ADT/StringRef.h"
-#include <memory>
-
-// Forward-declare LLVM classes.
-namespace llvm {
-class LLVMContext;
-class Module;
-} // namespace llvm
-
-namespace mlir {
-class Operation;
-
-/// Convert the given LLVM-module-like operation into NVVM IR. This conversion
-/// requires the registration of the LLVM IR dialect and will extract the LLVM
-/// context 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 `nullptr`.
-std::unique_ptr<llvm::Module>
-translateModuleToNVVMIR(Operation *m, llvm::LLVMContext &llvmContext,
-                        llvm::StringRef name = "LLVMDialectModule");
-
-} // namespace mlir
-
-#endif // MLIR_TARGET_NVVMIR_H

mlir/include/mlir/Target/ROCDLIR.h

-//===- ROCDLIR.h - MLIR to LLVM + ROCDL IR conversion -----------*- 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 declares the entry point for the MLIR to LLVM + ROCDL IR
-// conversion.
-//
-//===----------------------------------------------------------------------===//
-
-#ifndef MLIR_TARGET_ROCDLIR_H
-#define MLIR_TARGET_ROCDLIR_H
-
-#include "llvm/ADT/StringRef.h"
-#include <memory>
-
-// Forward-declare LLVM classes.
-namespace llvm {
-class LLVMContext;
-class Module;
-} // namespace llvm
-
-namespace mlir {
-class Operation;
-
-/// Convert the given LLVM-module-like operation into ROCDL IR. This conversion
-/// requires the registration of the LLVM IR dialect and will extract the LLVM
-/// context 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 `nullptr`.
-std::unique_ptr<llvm::Module>
-translateModuleToROCDLIR(Operation *m, llvm::LLVMContext &llvmContext,
-                         llvm::StringRef name = "LLVMDialectModule");
-
-} // namespace mlir
-
-#endif // MLIR_TARGET_ROCDLIR_H

mlir/lib/Target/CMakeLists.txt

add_mlir_translation_library(MLIRTargetLLVMIRModuleTranslation
 
   LINK_LIBS PUBLIC
   MLIRLLVMIR
   MLIROpenMP
   MLIRLLVMIRTransforms
   MLIRTranslation
   )
 
-add_mlir_translation_library(MLIRTargetAVX512
-  LLVMIR/LLVMAVX512Intr.cpp
-
-  ADDITIONAL_HEADER_DIRS
-  ${MLIR_MAIN_INCLUDE_DIR}/mlir/Target/LLVMIR
-
-  DEPENDS
-  MLIRLLVMAVX512ConversionsIncGen
-
-  LINK_COMPONENTS
-  Core
-
-  LINK_LIBS PUBLIC
-  MLIRIR
-  MLIRLLVMAVX512
-  MLIRLLVMIR
-  MLIRTargetLLVMIR
-  MLIRTargetLLVMIRModuleTranslation
-  )
-
 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
-  MLIRTargetLLVMIRModuleTranslation
-  )
-
-add_mlir_translation_library(MLIRTargetArmNeon
-  LLVMIR/LLVMArmNeonIntr.cpp
-
-  ADDITIONAL_HEADER_DIRS
-  ${MLIR_MAIN_INCLUDE_DIR}/mlir/Target/LLVMIR
-
-  DEPENDS
-  MLIRLLVMArmNeonConversionsIncGen
-
-  LINK_COMPONENTS
-  Core
-
-  LINK_LIBS PUBLIC
-  MLIRIR
-  MLIRLLVMArmNeon
-  MLIRLLVMIR
-  MLIRTargetLLVMIR
-  MLIRTargetLLVMIRModuleTranslation
-  )
-
-add_mlir_translation_library(MLIRTargetArmSVE
-  LLVMIR/LLVMArmSVEIntr.cpp
-
-  ADDITIONAL_HEADER_DIRS
-  ${MLIR_MAIN_INCLUDE_DIR}/mlir/Target/LLVMIR
-
-  DEPENDS
-  MLIRLLVMArmSVEConversionsIncGen
-
-  LINK_COMPONENTS
-  Core
-
-  LINK_LIBS PUBLIC
-  MLIRIR
-  MLIRLLVMArmSVE
-  MLIRLLVMIR
-  MLIRTargetLLVMIR
-  MLIRTargetLLVMIRModuleTranslation
-  )
-
-add_mlir_translation_library(MLIRTargetNVVMIR
-  LLVMIR/ConvertToNVVMIR.cpp
-
-  ADDITIONAL_HEADER_DIRS
-  ${MLIR_MAIN_INCLUDE_DIR}/mlir/Target/LLVMIR
-
-  DEPENDS
-  intrinsics_gen
-
-  LINK_COMPONENTS
-  Core
-
-  LINK_LIBS PUBLIC
-  MLIRGPU
-  MLIRIR
-  MLIRLLVMIR
-  MLIRNVVMIR
-  MLIRTargetLLVMIR
-  MLIRTargetLLVMIRModuleTranslation
-  )
-
-add_mlir_translation_library(MLIRTargetROCDLIR
-  LLVMIR/ConvertToROCDLIR.cpp
-
-  ADDITIONAL_HEADER_DIRS
-  ${MLIR_MAIN_INCLUDE_DIR}/mlir/Target/LLVMIR
-
-  DEPENDS
-  intrinsics_gen
-
-  LINK_COMPONENTS
-  Core
-
-  LINK_LIBS PUBLIC
-  MLIRGPU
-  MLIRIR
-  MLIRLLVMIR
-  MLIRROCDLIR
-  MLIRTargetLLVMIR
+  MLIRROCDLToLLVMIRTranslation
   MLIRTargetLLVMIRModuleTranslation
   )

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/Translation.h"
 
 #include "llvm/ADT/StringRef.h"
 #include "llvm/IR/Module.h"
 #include "llvm/IR/Verifier.h"
 #include "llvm/Support/ToolOutputFile.h"
 
 using namespace mlir;
 
 std::unique_ptr<llvm::Module>
-mlir::translateModuleToLLVMIR(ModuleOp m, llvm::LLVMContext &llvmContext,
+mlir::translateModuleToLLVMIR(Operation *op, llvm::LLVMContext &llvmContext,
                               StringRef name) {
   auto llvmModule =
-      LLVM::ModuleTranslation::translateModule<>(m, llvmContext, name);
+      LLVM::ModuleTranslation::translateModule<>(op, llvmContext, name);
   if (!llvmModule)
-    emitError(m.getLoc(), "Fail to convert MLIR to LLVM IR");
+    emitError(op->getLoc(), "Fail to convert MLIR to LLVM IR");
   else if (verifyModule(*llvmModule))
-    emitError(m.getLoc(), "LLVM IR fails to verify");
+    emitError(op->getLoc(), "LLVM IR fails to verify");
   return llvmModule;
 }
 
 void mlir::registerLLVMDialectTranslation(DialectRegistry &registry) {
   registry.insert<LLVM::LLVMDialect>();
   registry.addDialectInterface<LLVM::LLVMDialect,
                                LLVMDialectLLVMIRTranslationInterface>();
 }
TranslateFromMLIRRegistration registration(
             module, llvmContext, "LLVMDialectModule");
         if (!llvmModule)
           return failure();
 
         llvmModule->print(output, nullptr);
         return success();
       },
       [](DialectRegistry &registry) {
-        registry.insert<omp::OpenMPDialect>();
+        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);
       });
 }
 } // namespace mlir

mlir/lib/Target/LLVMIR/ConvertToNVVMIR.cpp

-//===- ConvertToNVVMIR.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 + NVVM dialects and
-// LLVM IR with NVVM intrinsics and metadata.
-//
-//===----------------------------------------------------------------------===//
-
-#include "mlir/Target/NVVMIR.h"
-
-#include "mlir/Dialect/GPU/GPUDialect.h"
-#include "mlir/Dialect/LLVMIR/LLVMDialect.h"
-#include "mlir/Dialect/LLVMIR/NVVMDialect.h"
-#include "mlir/IR/BuiltinOps.h"
-#include "mlir/Target/LLVMIR.h"
-#include "mlir/Target/LLVMIR/ModuleTranslation.h"
-#include "mlir/Translation.h"
-
-#include "llvm/ADT/StringRef.h"
-#include "llvm/IR/IntrinsicsNVPTX.h"
-#include "llvm/IR/Module.h"
-#include "llvm/Support/ToolOutputFile.h"
-
-using namespace mlir;
-
-static llvm::Value *createIntrinsicCall(llvm::IRBuilder<> &builder,
-                                        llvm::Intrinsic::ID intrinsic,
-                                        ArrayRef<llvm::Value *> args = {}) {
-  llvm::Module *module = builder.GetInsertBlock()->getModule();
-  llvm::Function *fn = llvm::Intrinsic::getDeclaration(module, intrinsic);
-  return builder.CreateCall(fn, args);
-}
-
-static llvm::Intrinsic::ID getShflBflyIntrinsicId(llvm::Type *resultType,
-                                                  bool withPredicate) {
-  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;
-}
-
-namespace {
-class ModuleTranslation : public LLVM::ModuleTranslation {
-public:
-  using LLVM::ModuleTranslation::ModuleTranslation;
-
-protected:
-  LogicalResult convertOperation(Operation &opInst,
-                                 llvm::IRBuilder<> &builder) override {
-
-#include "mlir/Dialect/LLVMIR/NVVMConversions.inc"
-
-    return LLVM::ModuleTranslation::convertOperation(opInst, builder);
-  }
-
-  /// Allow access to the constructor.
-  friend LLVM::ModuleTranslation;
-};
-} // namespace
-
-std::unique_ptr<llvm::Module>
-mlir::translateModuleToNVVMIR(Operation *m, llvm::LLVMContext &llvmContext,
-                              StringRef name) {
-  // Register the translation to LLVM IR if nobody else did before. This may
-  // happen if this translation is called inside a pass pipeline that converts
-  // GPU dialects to binary blobs without translating the rest of the code.
-  registerLLVMDialectTranslation(*m->getContext());
-
-  auto llvmModule = LLVM::ModuleTranslation::translateModule<ModuleTranslation>(
-      m, llvmContext, name);
-  if (!llvmModule)
-    return llvmModule;
-
-  // Insert the nvvm.annotations kernel so that the NVVM backend recognizes the
-  // function as a kernel.
-  for (auto func :
-       ModuleTranslation::getModuleBody(m).getOps<LLVM::LLVMFuncOp>()) {
-    if (!func->getAttrOfType<UnitAttr>(
-            NVVM::NVVMDialect::getKernelFuncAttrName()))
-      continue;
-
-    auto *llvmFunc = llvmModule->getFunction(func.getName());
-
-    llvm::Metadata *llvmMetadata[] = {
-        llvm::ValueAsMetadata::get(llvmFunc),
-        llvm::MDString::get(llvmModule->getContext(), "kernel"),
-        llvm::ValueAsMetadata::get(llvm::ConstantInt::get(
-            llvm::Type::getInt32Ty(llvmModule->getContext()), 1))};
-    llvm::MDNode *llvmMetadataNode =
-        llvm::MDNode::get(llvmModule->getContext(), llvmMetadata);
-    llvmModule->getOrInsertNamedMetadata("nvvm.annotations")
-        ->addOperand(llvmMetadataNode);
-  }
-
-  return llvmModule;
-}
-
-namespace mlir {
-void registerToNVVMIRTranslation() {
-  TranslateFromMLIRRegistration registration(
-      "mlir-to-nvvmir",
-      [](ModuleOp module, raw_ostream &output) {
-        llvm::LLVMContext llvmContext;
-        auto llvmModule = mlir::translateModuleToNVVMIR(module, llvmContext);
-        if (!llvmModule)
-          return failure();
-
-        llvmModule->print(output, nullptr);
-        return success();
-      },
-      [](DialectRegistry &registry) {
-        registry.insert<NVVM::NVVMDialect>();
-        registerLLVMDialectTranslation(registry);
-      });
-}
-} // namespace mlir

mlir/lib/Target/LLVMIR/ConvertToROCDLIR.cpp

-//===- ConvertToROCDLIR.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 + ROCDL dialects and
-// LLVM IR with ROCDL intrinsics and metadata.
-//
-//===----------------------------------------------------------------------===//
-
-#include "mlir/Target/ROCDLIR.h"
-
-#include "mlir/Dialect/GPU/GPUDialect.h"
-#include "mlir/Dialect/LLVMIR/LLVMDialect.h"
-#include "mlir/Dialect/LLVMIR/ROCDLDialect.h"
-#include "mlir/IR/BuiltinOps.h"
-#include "mlir/Target/LLVMIR.h"
-#include "mlir/Target/LLVMIR/ModuleTranslation.h"
-#include "mlir/Translation.h"
-
-#include "llvm/ADT/StringRef.h"
-#include "llvm/IR/IntrinsicsAMDGPU.h"
-#include "llvm/IR/Module.h"
-#include "llvm/Support/ToolOutputFile.h"
-
-using namespace mlir;
-
-// Create a call to llvm intrinsic
-static llvm::Value *createIntrinsicCall(llvm::IRBuilder<> &builder,
-                                        llvm::Intrinsic::ID intrinsic,
-                                        ArrayRef<llvm::Value *> args = {},
-                                        ArrayRef<llvm::Type *> tys = {}) {
-  llvm::Module *module = builder.GetInsertBlock()->getModule();
-  llvm::Function *fn = llvm::Intrinsic::getDeclaration(module, intrinsic, tys);
-  return builder.CreateCall(fn, args);
-}
-
-// Create a call to ROCm-Device-Library function
-//   Currently this routine will work only for calling ROCDL functions that
-// take a single int32 argument. It is likely that the interface of this
-// function will change to make it more generic.
-static llvm::Value *createDeviceFunctionCall(llvm::IRBuilder<> &builder,
-                                             StringRef fn_name, int parameter) {
-  llvm::Module *module = builder.GetInsertBlock()->getModule();
-  llvm::FunctionType *function_type = llvm::FunctionType::get(
-      llvm::Type::getInt64Ty(module->getContext()), // return type.
-      llvm::Type::getInt32Ty(module->getContext()), // parameter type.
-      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));
-}
-
-namespace {
-class ModuleTranslation : public LLVM::ModuleTranslation {
-public:
-  using LLVM::ModuleTranslation::ModuleTranslation;
-
-protected:
-  LogicalResult convertOperation(Operation &opInst,
-                                 llvm::IRBuilder<> &builder) override {
-
-#include "mlir/Dialect/LLVMIR/ROCDLConversions.inc"
-
-    return LLVM::ModuleTranslation::convertOperation(opInst, builder);
-  }
-
-  /// Allow access to the constructor.
-  friend LLVM::ModuleTranslation;
-};
-} // namespace
-
-std::unique_ptr<llvm::Module>
-mlir::translateModuleToROCDLIR(Operation *m, llvm::LLVMContext &llvmContext,
-                               StringRef name) {
-  // Register the translation to LLVM IR if nobody else did before. This may
-  // happen if this translation is called inside a pass pipeline that converts
-  // GPU dialects to binary blobs without translating the rest of the code.
-  registerLLVMDialectTranslation(*m->getContext());
-
-  // Lower MLIR (with RODL Dialect) to LLVM IR (with ROCDL intrinsics).
-  auto llvmModule = LLVM::ModuleTranslation::translateModule<ModuleTranslation>(
-      m, llvmContext, name);
-
-  // Foreach GPU kernel:
-  // 1. Insert AMDGPU_KERNEL calling convention.
-  // 2. Insert amdgpu-flat-workgroup-size(1, 1024) attribute.
-  for (auto func :
-       ModuleTranslation::getModuleBody(m).getOps<LLVM::LLVMFuncOp>()) {
-    if (!func->getAttrOfType<UnitAttr>(
-            ROCDL::ROCDLDialect::getKernelFuncAttrName()))
-      continue;
-
-    auto *llvmFunc = llvmModule->getFunction(func.getName());
-
-    llvmFunc->setCallingConv(llvm::CallingConv::AMDGPU_KERNEL);
-
-    llvmFunc->addFnAttr("amdgpu-flat-work-group-size", "1, 1024");
-  }
-
-  return llvmModule;
-}
-
-namespace mlir {
-void registerToROCDLIRTranslation() {
-  TranslateFromMLIRRegistration registration(
-      "mlir-to-rocdlir",
-      [](ModuleOp module, raw_ostream &output) {
-        llvm::LLVMContext llvmContext;
-        auto llvmModule = mlir::translateModuleToROCDLIR(module, llvmContext);
-        if (!llvmModule)
-          return failure();
-
-        llvmModule->print(output, nullptr);
-        return success();
-      },
-      [](DialectRegistry &registry) {
-        registry.insert<ROCDL::ROCDLDialect>();
-        registerLLVMDialectTranslation(registry);
-      });
-}
-} // namespace mlir

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

+add_subdirectory(LLVMArmNeon)
+add_subdirectory(LLVMArmSVE)
+add_subdirectory(LLVMAVX512)
 add_subdirectory(LLVMIR)
+add_subdirectory(NVVM)
 add_subdirectory(OpenMP)
+add_subdirectory(ROCDL)

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
+  )

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

+//===- LLVMAVX512ToLLVMIRTranslation.cpp - Translate LLVMAVX512 to LLVM IR-===//
+//
+// 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 LLVMAVX512 dialect and
+// LLVM IR.
+//
+//===----------------------------------------------------------------------===//
+
+#include "mlir/Target/LLVMIR/Dialect/LLVMAVX512/LLVMAVX512ToLLVMIRTranslation.h"
+#include "mlir/Dialect/LLVMIR/LLVMAVX512Dialect.h"
+#include "mlir/IR/Operation.h"
+#include "mlir/Target/LLVMIR/ModuleTranslation.h"
+
+#include "llvm/IR/IRBuilder.h"
+#include "llvm/IR/IntrinsicsX86.h"
+
+using namespace mlir;
+using namespace mlir::LLVM;
+
+LogicalResult
+mlir::LLVMAVX512DialectLLVMIRTranslationInterface::convertOperation(
+    Operation *op, llvm::IRBuilderBase &builder,
+    LLVM::ModuleTranslation &moduleTranslation) const {
+  Operation &opInst = *op;
+#include "mlir/Dialect/LLVMIR/LLVMAVX512Conversions.inc"
+
+  return failure();
+}

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
+  )

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

+//===- LLVMArmNeonToLLVMIRTranslation.cpp - LLVMArmNeon to LLVM IR --------===//
+//
+// 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 LLVMArmNeon dialect and
+// LLVM IR.
+//
+//===----------------------------------------------------------------------===//
+
+#include "mlir/Target/LLVMIR/Dialect/LLVMArmNeon/LLVMArmNeonToLLVMIRTranslation.h"
+#include "mlir/Dialect/LLVMIR/LLVMArmNeonDialect.h"
+#include "mlir/IR/Operation.h"
+#include "mlir/Target/LLVMIR/ModuleTranslation.h"
+
+#include "llvm/IR/IRBuilder.h"
+#include "llvm/IR/IntrinsicsAArch64.h"
+
+using namespace mlir;
+using namespace mlir::LLVM;
+
+LogicalResult
+mlir::LLVMArmNeonDialectLLVMIRTranslationInterface::convertOperation(
+    Operation *op, llvm::IRBuilderBase &builder,
+    LLVM::ModuleTranslation &moduleTranslation) const {
+  Operation &opInst = *op;
+#include "mlir/Dialect/LLVMIR/LLVMArmNeonConversions.inc"
+
+  return failure();
+}

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
+  )

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

+//===- LLVMArmSVEToLLVMIRTranslation.cpp - Translate LLVMArmSVE to LLVM IR-===//
+//
+// 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 LLVMArmSVE dialect and
+// LLVM IR.
+//
+//===----------------------------------------------------------------------===//
+
+#include "mlir/Target/LLVMIR/Dialect/LLVMArmSVE/LLVMArmSVEToLLVMIRTranslation.h"
+#include "mlir/Dialect/LLVMIR/LLVMArmSVEDialect.h"
+#include "mlir/IR/Operation.h"
+#include "mlir/Target/LLVMIR/ModuleTranslation.h"
+
+#include "llvm/IR/IRBuilder.h"
+#include "llvm/IR/IntrinsicsAArch64.h"
+
+using namespace mlir;
+using namespace mlir::LLVM;
+
+LogicalResult
+mlir::LLVMArmSVEDialectLLVMIRTranslationInterface::convertOperation(
+    Operation *op, llvm::IRBuilderBase &builder,
+    LLVM::ModuleTranslation &moduleTranslation) const {
+  Operation &opInst = *op;
+#include "mlir/Dialect/LLVMIR/LLVMArmSVEConversions.inc"
+
+  return failure();
+}

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

 
 #include "llvm/IR/IRBuilder.h"
 #include "llvm/IR/InlineAsm.h"
 #include "llvm/IR/MDBuilder.h"
 #include "llvm/IR/Operator.h"
 
 using namespace mlir;
 using namespace mlir::LLVM;
+using mlir::LLVM::detail::getLLVMConstant;
 
 #include "mlir/Dialect/LLVMIR/LLVMConversionEnumsToLLVM.inc"
 
 /// Convert MLIR integer comparison predicate to LLVM IR comparison predicate.
 static llvm::CmpInst::Predicate getLLVMCmpPredicate(ICmpPredicate p) {
   switch (p) {
   case LLVM::ICmpPredicate::eq:
     return llvm::CmpInst::Predicate::ICMP_EQ;
static llvm::FastMathFlags getFastmathFlags(FastmathFlagsInterface &op) {
   llvm::FastMathFlags ret;
   auto fmf = op.fastmathFlags();
   for (auto it : handlers)
     if (bitEnumContains(fmf, it.first))
       (ret.*(it.second))(true);
   return ret;
 }
 
-namespace {
-/// Dispatcher functional object targeting different overloads of
-/// ModuleTranslation::mapValue.
-// TODO: this is only necessary for compatibility with the code emitted from
-// ODS, remove when ODS is updated (after all dialects have migrated to the new
-// translation mechanism).
-struct MapValueDispatcher {
-  explicit MapValueDispatcher(ModuleTranslation &mt) : moduleTranslation(mt) {}
-
-  llvm::Value *&operator()(mlir::Value v) {
-    return moduleTranslation.mapValue(v);
-  }
-
-  void operator()(mlir::Value m, llvm::Value *l) {
-    moduleTranslation.mapValue(m, l);
-  }
-
-  LLVM::ModuleTranslation &moduleTranslation;
-};
-} // end namespace
-
 static LogicalResult
 convertOperationImpl(Operation &opInst, llvm::IRBuilderBase &builder,
                      LLVM::ModuleTranslation &moduleTranslation) {
   auto extractPosition = [](ArrayAttr attr) {
     SmallVector<unsigned, 4> position;
     position.reserve(attr.size());
     for (Attribute v : attr)
       position.push_back(v.cast<IntegerAttr>().getValue().getZExtValue());
     return position;
   };
 
   llvm::IRBuilder<>::FastMathFlagGuard fmfGuard(builder);
   if (auto fmf = dyn_cast<FastmathFlagsInterface>(opInst))
     builder.setFastMathFlags(getFastmathFlags(fmf));
 
-  // TODO: these are necessary for compatibility with the code emitted from ODS,
-  // remove them when ODS is updated (after all dialects have migrated to the
-  // new translation mechanism).
-  MapValueDispatcher mapValue(moduleTranslation);
-  auto lookupValue = [&](mlir::Value v) {
-    return moduleTranslation.lookupValue(v);
-  };
-  auto convertType = [&](Type ty) { return moduleTranslation.convertType(ty); };
-  auto lookupValues = [&](ValueRange vs) {
-    return moduleTranslation.lookupValues(vs);
-  };
-  auto getLLVMConstant = [&](llvm::Type *ty, Attribute attr, Location loc) {
-    return moduleTranslation.getLLVMConstant(ty, attr, loc);
-  };
-
 #include "mlir/Dialect/LLVMIR/LLVMConversions.inc"
 
   // Emit function calls.  If the "callee" attribute is present, this is a
   // direct function call and we also need to look up the remapped function
   // itself.  Otherwise, this is an indirect call and the callee is the first
   // operand, look it up as a normal value.  Return the llvm::Value representing
   // the function result, which may be of llvm::VoidTy type.
   auto convertCall = [&](Operation &op) -> llvm::Value * {
return builder.CreateCall(calleeType, operandsRef.front(),
                               operandsRef.drop_front());
   };
 
   // Emit calls.  If the called function has a result, remap the corresponding
   // value.  Note that LLVM IR dialect CallOp has either 0 or 1 result.
   if (isa<LLVM::CallOp>(opInst)) {
     llvm::Value *result = convertCall(opInst);
     if (opInst.getNumResults() != 0) {
-      mapValue(opInst.getResult(0), result);
+      moduleTranslation.mapValue(opInst.getResult(0), result);
       return success();
     }
     // Check that LLVM call returns void for 0-result functions.
     return success(result->getType()->isVoidTy());
   }
 
   if (auto inlineAsmOp = dyn_cast<LLVM::InlineAsmOp>(opInst)) {
     // TODO: refactor function type creation which usually occurs in std-LLVM
if (auto inlineAsmOp = dyn_cast<LLVM::InlineAsmOp>(opInst)) {
     } else {
       assert(inlineAsmOp.getNumResults() == 1);
       resultType = inlineAsmOp.getResultTypes()[0];
     }
     auto ft = LLVM::LLVMFunctionType::get(resultType, operandTypes);
     llvm::InlineAsm *inlineAsmInst =
         inlineAsmOp.asm_dialect().hasValue()
             ? llvm::InlineAsm::get(
-                  static_cast<llvm::FunctionType *>(convertType(ft)),
+                  static_cast<llvm::FunctionType *>(
+                      moduleTranslation.convertType(ft)),
                   inlineAsmOp.asm_string(), inlineAsmOp.constraints(),
                   inlineAsmOp.has_side_effects(), inlineAsmOp.is_align_stack(),
                   convertAsmDialectToLLVM(*inlineAsmOp.asm_dialect()))
             : llvm::InlineAsm::get(
-                  static_cast<llvm::FunctionType *>(convertType(ft)),
+                  static_cast<llvm::FunctionType *>(
+                      moduleTranslation.convertType(ft)),
                   inlineAsmOp.asm_string(), inlineAsmOp.constraints(),
                   inlineAsmOp.has_side_effects(), inlineAsmOp.is_align_stack());
-    llvm::Value *result =
-        builder.CreateCall(inlineAsmInst, lookupValues(inlineAsmOp.operands()));
+    llvm::Value *result = builder.CreateCall(
+        inlineAsmInst, moduleTranslation.lookupValues(inlineAsmOp.operands()));
     if (opInst.getNumResults() != 0)
-      mapValue(opInst.getResult(0), result);
+      moduleTranslation.mapValue(opInst.getResult(0), result);
     return success();
   }
 
   if (auto invOp = dyn_cast<LLVM::InvokeOp>(opInst)) {
-    auto operands = lookupValues(opInst.getOperands());
+    auto operands = moduleTranslation.lookupValues(opInst.getOperands());
     ArrayRef<llvm::Value *> operandsRef(operands);
     if (auto attr = opInst.getAttrOfType<FlatSymbolRefAttr>("callee")) {
       builder.CreateInvoke(moduleTranslation.lookupFunction(attr.getValue()),
                            moduleTranslation.lookupBlock(invOp.getSuccessor(0)),
                            moduleTranslation.lookupBlock(invOp.getSuccessor(1)),
                            operandsRef);
     } else {
       auto *calleePtrType =
           cast<llvm::PointerType>(operandsRef.front()->getType());
       auto *calleeType =
           cast<llvm::FunctionType>(calleePtrType->getElementType());
       builder.CreateInvoke(calleeType, operandsRef.front(),
                            moduleTranslation.lookupBlock(invOp.getSuccessor(0)),
                            moduleTranslation.lookupBlock(invOp.getSuccessor(1)),
                            operandsRef.drop_front());
     }
     return success();
   }
 
   if (auto lpOp = dyn_cast<LLVM::LandingpadOp>(opInst)) {
-    llvm::Type *ty = convertType(lpOp.getType());
+    llvm::Type *ty = moduleTranslation.convertType(lpOp.getType());
     llvm::LandingPadInst *lpi =
         builder.CreateLandingPad(ty, lpOp.getNumOperands());
 
     // Add clauses
-    for (llvm::Value *operand : lookupValues(lpOp.getOperands())) {
+    for (llvm::Value *operand :
+         moduleTranslation.lookupValues(lpOp.getOperands())) {
       // All operands should be constant - checked by verifier
       if (auto *constOperand = dyn_cast<llvm::Constant>(operand))
         lpi->addClause(constOperand);
     }
-    mapValue(lpOp.getResult(), lpi);
+    moduleTranslation.mapValue(lpOp.getResult(), lpi);
     return success();
   }
 
   // Emit branches.  We need to look up the remapped blocks and ignore the block
   // arguments that were transformed into PHI nodes.
   if (auto brOp = dyn_cast<LLVM::BrOp>(opInst)) {
     llvm::BranchInst *branch =
         builder.CreateBr(moduleTranslation.lookupBlock(brOp.getSuccessor()));
if (auto weights = switchOp.branch_weights()) {
                           .createBranchWeights(weightValues);
     }
 
     llvm::SwitchInst *switchInst = builder.CreateSwitch(
         moduleTranslation.lookupValue(switchOp.value()),
         moduleTranslation.lookupBlock(switchOp.defaultDestination()),
         switchOp.caseDestinations().size(), branchWeights);
 
-    auto *ty =
-        llvm::cast<llvm::IntegerType>(convertType(switchOp.value().getType()));
+    auto *ty = llvm::cast<llvm::IntegerType>(
+        moduleTranslation.convertType(switchOp.value().getType()));
     for (auto i :
          llvm::zip(switchOp.case_values()->cast<DenseIntElementsAttr>(),
                    switchOp.caseDestinations()))
       switchInst->addCase(
           llvm::ConstantInt::get(ty, std::get<0>(i).getLimitedValue()),
           moduleTranslation.lookupBlock(std::get<1>(i)));
 
     moduleTranslation.mapBranch(&opInst, switchInst);
     return success();
   }
 
   // Emit addressof.  We need to look up the global value referenced by the
   // operation and store it in the MLIR-to-LLVM value mapping.  This does not
   // emit any LLVM instruction.
   if (auto addressOfOp = dyn_cast<LLVM::AddressOfOp>(opInst)) {
     LLVM::GlobalOp global = addressOfOp.getGlobal();
     LLVM::LLVMFuncOp function = addressOfOp.getFunction();
 
     // The verifier should not have allowed this.
     assert((global || function) &&
            "referencing an undefined global or function");
 
-    mapValue(addressOfOp.getResult(),
-             global ? moduleTranslation.lookupGlobal(global)
-                    : moduleTranslation.lookupFunction(function.getName()));
+    moduleTranslation.mapValue(
+        addressOfOp.getResult(),
     return success();
   }
 
   return failure();
 }
 
 LogicalResult mlir::LLVMDialectLLVMIRTranslationInterface::convertOperation(
     Operation *op, llvm::IRBuilderBase &builder,
     LLVM::ModuleTranslation &moduleTranslation) const {
   return convertOperationImpl(*op, builder, moduleTranslation);
 }

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
+  )

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

+//===- NVVMToLLVMIRTranslation.cpp - Translate NVVM to LLVM IR ------------===//
+//
+// 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 NVVM dialect and
+// LLVM IR.
+//
+//===----------------------------------------------------------------------===//
+
+#include "mlir/Target/LLVMIR/Dialect/NVVM/NVVMToLLVMIRTranslation.h"
+#include "mlir/Dialect/LLVMIR/NVVMDialect.h"
+#include "mlir/IR/Operation.h"
+#include "mlir/Target/LLVMIR/ModuleTranslation.h"
+
+#include "llvm/IR/IRBuilder.h"
+#include "llvm/IR/IntrinsicsNVPTX.h"
+
+using namespace mlir;
+using namespace mlir::LLVM;
+using mlir::LLVM::detail::createIntrinsicCall;
+
+static llvm::Intrinsic::ID getShflBflyIntrinsicId(llvm::Type *resultType,
+                                                  bool withPredicate) {
+  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 {
+  Operation &opInst = *op;
+#include "mlir/Dialect/LLVMIR/NVVMConversions.inc"
+
+  return failure();
+}
+
+LogicalResult mlir::NVVMDialectLLVMIRTranslationInterface::amendOperation(
+    Operation *op, NamedAttribute attribute,
+    LLVM::ModuleTranslation &moduleTranslation) const {
+  if (attribute.first == NVVM::NVVMDialect::getKernelFuncAttrName()) {
+    auto func = dyn_cast<LLVM::LLVMFuncOp>(op);
+    if (!func)
+      return failure();
+
+    llvm::LLVMContext &llvmContext = moduleTranslation.getLLVMContext();
+    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();
+}

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
+  )

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

+//===- ROCDLToLLVMIRTranslation.cpp - Translate ROCDL to LLVM IR ----------===//
+//
+// 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 ROCDL dialect and
+// LLVM IR.
+//
+//===----------------------------------------------------------------------===//
+
+#include "mlir/Target/LLVMIR/Dialect/ROCDL/ROCDLToLLVMIRTranslation.h"
+#include "mlir/Dialect/LLVMIR/ROCDLDialect.h"
+#include "mlir/IR/Operation.h"
+#include "mlir/Target/LLVMIR/ModuleTranslation.h"
+
+#include "llvm/IR/IRBuilder.h"
+#include "llvm/IR/IntrinsicsAMDGPU.h"
+
+using namespace mlir;
+using namespace mlir::LLVM;
+using mlir::LLVM::detail::createIntrinsicCall;
+
+// Create a call to ROCm-Device-Library function
+// Currently this routine will work only for calling ROCDL functions that
+// take a single int32 argument. It is likely that the interface of this
+// function will change to make it more generic.
+static llvm::Value *createDeviceFunctionCall(llvm::IRBuilderBase &builder,
+                                             StringRef fn_name, int parameter) {
+  llvm::Module *module = builder.GetInsertBlock()->getModule();
+  llvm::FunctionType *function_type = llvm::FunctionType::get(
+      llvm::Type::getInt64Ty(module->getContext()), // return type.
+      llvm::Type::getInt32Ty(module->getContext()), // parameter type.
+      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 {
+  Operation &opInst = *op;
+#include "mlir/Dialect/LLVMIR/ROCDLConversions.inc"
+
+  return failure();
+}
+
+LogicalResult mlir::ROCDLDialectLLVMIRTranslationInterface::amendOperation(
+    Operation *op, NamedAttribute attribute,
+    LLVM::ModuleTranslation &moduleTranslation) const {
+  if (attribute.first == ROCDL::ROCDLDialect::getKernelFuncAttrName()) {
+    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());
+    llvmFunc->setCallingConv(llvm::CallingConv::AMDGPU_KERNEL);
+    llvmFunc->addFnAttr("amdgpu-flat-work-group-size", "1, 1024");
+  }
+  return success();
+}

mlir/lib/Target/LLVMIR/LLVMAVX512Intr.cpp

-//===- AVX512Intr.cpp - Convert MLIR LLVM dialect to LLVM intrinsics ------===//
-//
-// 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 and AVX512 dialects
-// and LLVM IR with AVX intrinsics.
-//
-//===----------------------------------------------------------------------===//
-
-#include "mlir/Dialect/LLVMIR/LLVMAVX512Dialect.h"
-#include "mlir/Target/LLVMIR.h"
-#include "mlir/Target/LLVMIR/ModuleTranslation.h"
-#include "mlir/Translation.h"
-#include "llvm/IR/IntrinsicsX86.h"
-
-using namespace mlir;
-
-namespace {
-class LLVMAVX512ModuleTranslation : public LLVM::ModuleTranslation {
-  friend LLVM::ModuleTranslation;
-
-public:
-  using LLVM::ModuleTranslation::ModuleTranslation;
-
-protected:
-  LogicalResult convertOperation(Operation &opInst,
-                                 llvm::IRBuilder<> &builder) override {
-#include "mlir/Dialect/LLVMIR/LLVMAVX512Conversions.inc"
-
-    return LLVM::ModuleTranslation::convertOperation(opInst, builder);
-  }
-};
-
-std::unique_ptr<llvm::Module>
-translateLLVMAVX512ModuleToLLVMIR(Operation *m, llvm::LLVMContext &llvmContext,
-                                  StringRef name) {
-  return LLVM::ModuleTranslation::translateModule<LLVMAVX512ModuleTranslation>(
-      m, llvmContext, name);
-}
-} // end namespace
-
-namespace mlir {
-void registerAVX512ToLLVMIRTranslation() {
-  TranslateFromMLIRRegistration reg(
-      "avx512-mlir-to-llvmir",
-      [](ModuleOp module, raw_ostream &output) {
-        llvm::LLVMContext llvmContext;
-        auto llvmModule = translateLLVMAVX512ModuleToLLVMIR(
-            module, llvmContext, "LLVMDialectModule");
-        if (!llvmModule)
-          return failure();
-
-        llvmModule->print(output, nullptr);
-        return success();
-      },
-      [](DialectRegistry &registry) {
-        registry.insert<LLVM::LLVMAVX512Dialect>();
-        registerLLVMDialectTranslation(registry);
-      });
-}
-} // namespace mlir

mlir/lib/Target/LLVMIR/LLVMArmNeonIntr.cpp

-//===- ArmNeonIntr.cpp - Convert MLIR LLVM dialect to LLVM intrinsics -----===//
-//
-// 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 and ArmNeon dialects
-// and LLVM IR with ArmNeon intrinsics.
-//
-//===----------------------------------------------------------------------===//
-
-#include "mlir/Dialect/LLVMIR/LLVMArmNeonDialect.h"
-#include "mlir/Target/LLVMIR.h"
-#include "mlir/Target/LLVMIR/ModuleTranslation.h"
-#include "mlir/Translation.h"
-#include "llvm/IR/IntrinsicsAArch64.h"
-
-using namespace mlir;
-
-namespace {
-class LLVMArmNeonModuleTranslation : public LLVM::ModuleTranslation {
-  friend LLVM::ModuleTranslation;
-
-public:
-  using LLVM::ModuleTranslation::ModuleTranslation;
-
-protected:
-  LogicalResult convertOperation(Operation &opInst,
-                                 llvm::IRBuilder<> &builder) override {
-#include "mlir/Dialect/LLVMIR/LLVMArmNeonConversions.inc"
-
-    return LLVM::ModuleTranslation::convertOperation(opInst, builder);
-  }
-};
-
-std::unique_ptr<llvm::Module>
-translateLLVMArmNeonModuleToLLVMIR(Operation *m, llvm::LLVMContext &llvmContext,
-                                   StringRef name) {
-  return LLVM::ModuleTranslation::translateModule<LLVMArmNeonModuleTranslation>(
-      m, llvmContext, name);
-}
-} // end namespace
-
-namespace mlir {
-void registerArmNeonToLLVMIRTranslation() {
-  TranslateFromMLIRRegistration reg(
-      "arm-neon-mlir-to-llvmir",
-      [](ModuleOp module, raw_ostream &output) {
-        llvm::LLVMContext llvmContext;
-        auto llvmModule = translateLLVMArmNeonModuleToLLVMIR(
-            module, llvmContext, "LLVMDialectModule");
-        if (!llvmModule)
-          return failure();
-
-        llvmModule->print(output, nullptr);
-        return success();
-      },
-      [](DialectRegistry &registry) {
-        registry.insert<LLVM::LLVMArmNeonDialect>();
-        registerLLVMDialectTranslation(registry);
-      });
-}
-} // namespace mlir

mlir/lib/Target/LLVMIR/LLVMArmSVEIntr.cpp

-//===- LLVMArmSVEIntr.cpp - Convert MLIR LLVM dialect to LLVM intrinsics --===//
-//
-// 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 and ArmSVE dialects
-// and LLVM IR with Arm SVE intrinsics.
-//
-//===----------------------------------------------------------------------===//
-
-#include "mlir/Dialect/LLVMIR/LLVMArmSVEDialect.h"
-#include "mlir/Target/LLVMIR.h"
-#include "mlir/Target/LLVMIR/ModuleTranslation.h"
-#include "mlir/Translation.h"
-#include "llvm/IR/IntrinsicsAArch64.h"
-
-using namespace mlir;
-
-namespace {
-class LLVMArmSVEModuleTranslation : public LLVM::ModuleTranslation {
-  friend LLVM::ModuleTranslation;
-
-public:
-  using LLVM::ModuleTranslation::ModuleTranslation;
-
-protected:
-  LogicalResult convertOperation(Operation &opInst,
-                                 llvm::IRBuilder<> &builder) override {
-#include "mlir/Dialect/LLVMIR/LLVMArmSVEConversions.inc"
-
-    return LLVM::ModuleTranslation::convertOperation(opInst, builder);
-  }
-};
-} // end namespace
-
-static std::unique_ptr<llvm::Module>
-translateLLVMArmSVEModuleToLLVMIR(Operation *m, llvm::LLVMContext &llvmContext,
-                                  StringRef name) {
-  return LLVM::ModuleTranslation::translateModule<LLVMArmSVEModuleTranslation>(
-      m, llvmContext, name);
-}
-
-namespace mlir {
-void registerArmSVEToLLVMIRTranslation() {
-  TranslateFromMLIRRegistration reg(
-      "arm-sve-mlir-to-llvmir",
-      [](ModuleOp module, raw_ostream &output) {
-        llvm::LLVMContext llvmContext;
-        auto llvmModule = translateLLVMArmSVEModuleToLLVMIR(
-            module, llvmContext, "LLVMDialectModule");
-        if (!llvmModule)
-          return failure();
-
-        llvmModule->print(output, nullptr);
-        return success();
-      },
-      [](DialectRegistry &registry) {
-        registry.insert<LLVM::LLVMArmSVEDialect>();
-        registerLLVMDialectTranslation(registry);
-      });
-}
-} // namespace mlir

mlir/lib/Target/LLVMIR/ModuleTranslation.cpp

do {
     }
   } while (true);
 }
 
 /// Create an LLVM IR constant of `llvmType` from the MLIR attribute `attr`.
 /// This currently supports integer, floating point, splat and dense element
 /// attributes and combinations thereof.  In case of error, report it to `loc`
 /// and return nullptr.
-llvm::Constant *ModuleTranslation::getLLVMConstant(llvm::Type *llvmType,
-                                                   Attribute attr,
-                                                   Location loc) {
+llvm::Constant *mlir::LLVM::detail::getLLVMConstant(
+    llvm::Type *llvmType, Attribute attr, Location loc,
+    const ModuleTranslation &moduleTranslation) {
   if (!attr)
     return llvm::UndefValue::get(llvmType);
   if (llvmType->isStructTy()) {
     emitError(loc, "struct types are not supported in constants");
     return nullptr;
   }
   // For integer types, we allow a mismatch in sizes as the index type in
   // MLIR might have a different size than the index type in the LLVM module.
   if (auto intAttr = attr.dyn_cast<IntegerAttr>())
     return llvm::ConstantInt::get(
         llvmType,
         intAttr.getValue().sextOrTrunc(llvmType->getIntegerBitWidth()));
   if (auto floatAttr = attr.dyn_cast<FloatAttr>())
     return llvm::ConstantFP::get(llvmType, floatAttr.getValue());
   if (auto funcAttr = attr.dyn_cast<FlatSymbolRefAttr>())
-    return llvm::ConstantExpr::getBitCast(lookupFunction(funcAttr.getValue()),
-                                          llvmType);
+    return llvm::ConstantExpr::getBitCast(
+        moduleTranslation.lookupFunction(funcAttr.getValue()), llvmType);
   if (auto splatAttr = attr.dyn_cast<SplatElementsAttr>()) {
     llvm::Type *elementType;
     uint64_t numElements;
     if (auto *arrayTy = dyn_cast<llvm::ArrayType>(llvmType)) {
       elementType = arrayTy->getElementType();
       numElements = arrayTy->getNumElements();
     } else {
       auto *vectorTy = cast<llvm::FixedVectorType>(llvmType);
       elementType = vectorTy->getElementType();
       numElements = vectorTy->getNumElements();
     }
     // Splat value is a scalar. Extract it only if the element type is not
     // another sequence type. The recursion terminates because each step removes
     // one outer sequential type.
     bool elementTypeSequential =
         isa<llvm::ArrayType, llvm::VectorType>(elementType);
     llvm::Constant *child = getLLVMConstant(
         elementType,
-        elementTypeSequential ? splatAttr : splatAttr.getSplatValue(), loc);
+        elementTypeSequential ? splatAttr : splatAttr.getSplatValue(), loc,
+        moduleTranslation);
     if (!child)
       return nullptr;
     if (llvmType->isVectorTy())
       return llvm::ConstantVector::getSplat(
           llvm::ElementCount::get(numElements, /*Scalable=*/false), child);
     if (llvmType->isArrayTy()) {
       auto *arrayType = llvm::ArrayType::get(elementType, numElements);
       SmallVector<llvm::Constant *, 8> constants(numElements, child);
       return llvm::ConstantArray::get(arrayType, constants);
     }
   }
 
   if (auto elementsAttr = attr.dyn_cast<ElementsAttr>()) {
     assert(elementsAttr.getType().hasStaticShape());
     assert(elementsAttr.getNumElements() != 0 &&
            "unexpected empty elements attribute");
     assert(!elementsAttr.getType().getShape().empty() &&
            "unexpected empty elements attribute shape");
 
     SmallVector<llvm::Constant *, 8> constants;
     constants.reserve(elementsAttr.getNumElements());
     llvm::Type *innermostType = getInnermostElementType(llvmType);
     for (auto n : elementsAttr.getValues<Attribute>()) {
-      constants.push_back(getLLVMConstant(innermostType, n, loc));
+      constants.push_back(
+          getLLVMConstant(innermostType, n, loc, moduleTranslation));
       if (!constants.back())
         return nullptr;
     }
     ArrayRef<llvm::Constant *> constantsRef = constants;
     llvm::Constant *result = buildSequentialConstant(
         constantsRef, elementsAttr.getType().getShape(), llvmType, loc);
     assert(constantsRef.empty() && "did not consume all elemental constants");
     return result;
   }
 
   if (auto stringAttr = attr.dyn_cast<StringAttr>()) {
     return llvm::ConstantDataArray::get(
-        llvmModule->getContext(), ArrayRef<char>{stringAttr.getValue().data(),
-                                                 stringAttr.getValue().size()});
+        moduleTranslation.getLLVMContext(),
+        ArrayRef<char>{stringAttr.getValue().data(),
   }
   emitError(loc, "unsupported constant value");
   return nullptr;
 }
 
 ModuleTranslation::ModuleTranslation(Operation *module,
                                      std::unique_ptr<llvm::Module> llvmModule)
     : mlirModule(module), llvmModule(std::move(llvmModule)),
for (Block &b : region) {
     }
   }
   assert(blocks.size() == region.getBlocks().size() &&
          "some blocks are not sorted");
 
   return blocks;
 }
 
+llvm::Value *mlir::LLVM::detail::createIntrinsicCall(
+    llvm::IRBuilderBase &builder, llvm::Intrinsic::ID intrinsic,
+    ArrayRef<llvm::Value *> args, ArrayRef<llvm::Type *> tys) {
+  llvm::Module *module = builder.GetInsertBlock()->getModule();
+  llvm::Function *fn = llvm::Intrinsic::getDeclaration(module, intrinsic, tys);
+  return builder.CreateCall(fn, args);
+}
+
 /// Given a single MLIR operation, create the corresponding LLVM IR operation
-/// using the `builder`.  LLVM IR Builder does not have a generic interface so
-/// this has to be a long chain of `if`s calling different functions with a
-/// different number of arguments.
+/// using the `builder`.
 LogicalResult ModuleTranslation::convertOperation(Operation &opInst,
                                                   llvm::IRBuilder<> &builder) {
-  if (succeeded(iface.convertOperation(&opInst, builder, *this)))
-    return success();
-
+  if (failed(iface.convertOperation(&opInst, builder, *this)))
     return opInst.emitError("unsupported or non-LLVM operation: ")
            << opInst.getName();
+
+  return convertDialectAttributes(&opInst);
 }
 
 /// Convert block to LLVM IR.  Unless `ignoreArguments` is set, emit PHI nodes
 /// to define values corresponding to the MLIR block arguments.  These nodes
 /// are not connected to the source basic blocks, which may not exist yet.  Uses
 /// `builder` to construct the LLVM IR. Expects the LLVM IR basic block to have
 /// been created for `bb` and included in the block mapping.  Inserts new
 /// instructions at the end of the block and leaves `builder` in a state
for (auto op : getModuleBody(mlirModule).getOps<LLVM::GlobalOp>()) {
     llvm::Constant *cst = llvm::UndefValue::get(type);
     if (op.getValueOrNull()) {
       // String attributes are treated separately because they cannot appear as
       // in-function constants and are thus not supported by getLLVMConstant.
       if (auto strAttr = op.getValueOrNull().dyn_cast_or_null<StringAttr>()) {
         cst = llvm::ConstantDataArray::getString(
             llvmModule->getContext(), strAttr.getValue(), /*AddNull=*/false);
         type = cst->getType();
-      } else if (!(cst = getLLVMConstant(type, op.getValueOrNull(),
-                                         op.getLoc()))) {
+      } else if (!(cst = getLLVMConstant(type, op.getValueOrNull(), op.getLoc(),
+                                         *this))) {
         return failure();
       }
     } else if (Block *initializer = op.getInitializerBlock()) {
       llvm::IRBuilder<> builder(llvmModule->getContext());
       for (auto &op : initializer->without_terminator()) {
         if (failed(convertOperation(op, builder)) ||
             !isa<llvm::Constant>(lookupValue(op.getResult(0))))
           return emitError(op.getLoc(), "unemittable constant value");
for (auto kvp : llvm::zip(func.getArguments(), llvmFunc->args())) {
     mapValue(mlirArg, &llvmArg);
     argIdx++;
   }
 
   // Check the personality and set it.
   if (func.personality().hasValue()) {
     llvm::Type *ty = llvm::Type::getInt8PtrTy(llvmFunc->getContext());
     if (llvm::Constant *pfunc =
-            getLLVMConstant(ty, func.personalityAttr(), func.getLoc()))
+            getLLVMConstant(ty, func.personalityAttr(), func.getLoc(), *this))
       llvmFunc->setPersonalityFn(pfunc);
   }
 
   // First, create all blocks so we can jump to them.
   llvm::LLVMContext &llvmContext = llvmFunc->getContext();
   for (auto &bb : func) {
     auto *llvmBB = llvm::BasicBlock::Create(llvmContext);
     llvmBB->insertInto(llvmFunc);
     mapBlock(&bb, llvmBB);
   }
 
   // Then, convert blocks one by one in topological order to ensure defs are
   // converted before uses.
   auto blocks = detail::getTopologicallySortedBlocks(func.getBody());
   for (Block *bb : blocks) {
     llvm::IRBuilder<> builder(llvmContext);
     if (failed(convertBlock(*bb, bb->isEntryBlock(), builder)))
       return failure();
   }
 
-  // Finally, after all blocks have been traversed and values mapped, connect
-  // the PHI nodes to the results of preceding blocks.
+  // After all blocks have been traversed and values mapped, connect the PHI
+  // nodes to the results of preceding blocks.
   detail::connectPHINodes(func.getBody(), *this);
+
+  // Finally, convert dialect attributes attached to the function.
+  return convertDialectAttributes(func);
+}
+
+LogicalResult ModuleTranslation::convertDialectAttributes(Operation *op) {
+  for (NamedAttribute attribute : op->getDialectAttrs())
+    if (failed(iface.amendOperation(op, attribute, *this)))
+      return failure();
   return success();
 }
 
 LogicalResult ModuleTranslation::checkSupportedModuleOps(Operation *m) {
   for (Operation &o : getModuleBody(m).getOperations())
     if (!isa<LLVM::LLVMFuncOp, LLVM::GlobalOp>(&o) &&
         !o.hasTrait<OpTrait::IsTerminator>())
       return o.emitOpError("unsupported module-level operation");
ModuleTranslation::lookupValues(ValueRange values) {
   return remapped;
 }
 
 const llvm::DILocation *
 ModuleTranslation::translateLoc(Location loc, llvm::DILocalScope *scope) {
   return debugTranslation->translateLoc(loc, scope);
 }
 
+llvm::NamedMDNode *
+ModuleTranslation::getOrInsertNamedModuleMetadata(StringRef name) {
+  return llvmModule->getOrInsertNamedMetadata(name);
+}
+
 std::unique_ptr<llvm::Module> ModuleTranslation::prepareLLVMModule(
     Operation *m, llvm::LLVMContext &llvmContext, StringRef name) {
   m->getContext()->getOrLoadDialect<LLVM::LLVMDialect>();
   auto llvmModule = std::make_unique<llvm::Module>(name, llvmContext);
   if (auto dataLayoutAttr =
           m->getAttr(LLVM::LLVMDialect::getDataLayoutAttrName()))
     llvmModule->setDataLayout(dataLayoutAttr.cast<StringAttr>().getValue());
   if (auto targetTripleAttr =

mlir/test/Target/arm-neon.mlir

-// RUN: mlir-opt -verify-diagnostics %s | mlir-opt | mlir-translate -arm-neon-mlir-to-llvmir | FileCheck %s
+// RUN: mlir-translate -mlir-to-llvmir %s | FileCheck %s
 
 // CHECK-LABEL: arm_neon_smull
 llvm.func @arm_neon_smull(%arg0: vector<8xi8>, %arg1: vector<8xi8>) -> !llvm.struct<(vector<8xi16>, vector<4xi32>, vector<2xi64>)> {
   //      CHECK: %[[V0:.*]] = call <8 x i16> @llvm.aarch64.neon.smull.v8i16(<8 x i8> %{{.*}}, <8 x i8> %{{.*}})
   // CHECK-NEXT: %[[V00:.*]] = shufflevector <8 x i16> %3, <8 x i16> %[[V0]], <4 x i32> <i32 3, i32 4, i32 5, i32 6>
   %0 = "llvm_arm_neon.smull"(%arg0, %arg1) : (vector<8xi8>, vector<8xi8>) -> vector<8xi16>
   %1 = llvm.shufflevector %0, %0 [3, 4, 5, 6] : vector<8xi16>, vector<8xi16>
 

mlir/test/Target/arm-sve.mlir

-// RUN: mlir-opt -verify-diagnostics %s | mlir-opt | mlir-translate --arm-sve-mlir-to-llvmir | FileCheck %s
+// RUN: mlir-translate --mlir-to-llvmir %s | FileCheck %s
 
 // CHECK-LABEL: define <vscale x 4 x i32> @arm_sve_sdot
 llvm.func @arm_sve_sdot(%arg0: !llvm.vec<?x16 x i8>,
                         %arg1: !llvm.vec<?x16 x i8>,
                         %arg2: !llvm.vec<?x4 x i32>)
                         -> !llvm.vec<?x4 x i32> {
   // CHECK: call <vscale x 4 x i32> @llvm.aarch64.sve.sdot.nxv4i32(<vscale x 4
   %0 = "llvm_arm_sve.sdot"(%arg2, %arg0, %arg1) :

mlir/test/Target/avx512.mlir

-// RUN: mlir-opt -verify-diagnostics %s | mlir-opt | mlir-translate --avx512-mlir-to-llvmir | FileCheck %s
+// RUN: mlir-translate --mlir-to-llvmir %s | FileCheck %s
 
 // CHECK-LABEL: define <16 x float> @LLVM_x86_avx512_mask_ps_512
 llvm.func @LLVM_x86_avx512_mask_ps_512(%a: vector<16 x f32>,
                                        %b: i32,
                                        %c: i16)
   -> (vector<16 x f32>)
 {
   // CHECK: call <16 x float> @llvm.x86.avx512.mask.rndscale.ps.512(<16 x float>

mlir/test/Target/nvvmir.mlir

-// RUN: mlir-translate -mlir-to-nvvmir %s | FileCheck %s
+// RUN: mlir-translate -mlir-to-llvmir %s | FileCheck %s
 
 llvm.func @nvvm_special_regs() -> i32 {
   // CHECK: %1 = call i32 @llvm.nvvm.read.ptx.sreg.tid.x()
   %1 = nvvm.read.ptx.sreg.tid.x : i32
   // CHECK: call i32 @llvm.nvvm.read.ptx.sreg.tid.y()
   %2 = nvvm.read.ptx.sreg.tid.y : i32
   // CHECK: call i32 @llvm.nvvm.read.ptx.sreg.tid.z()
   %3 = nvvm.read.ptx.sreg.tid.z : i32

mlir/test/Target/rocdl.mlir

-// RUN: mlir-translate -mlir-to-rocdlir %s | FileCheck %s
+// RUN: mlir-translate -mlir-to-llvmir %s | FileCheck %s
 
 llvm.func @rocdl_special_regs() -> i32 {
   // CHECK-LABEL: rocdl_special_regs
   // CHECK: call i32 @llvm.amdgcn.workitem.id.x()
   %1 = rocdl.workitem.id.x : i32
   // CHECK: call i32 @llvm.amdgcn.workitem.id.y()
   %2 = rocdl.workitem.id.y : i32
   // CHECK: call i32 @llvm.amdgcn.workitem.id.z()

mlir/test/lib/Transforms/CMakeLists.txt

add_mlir_library(MLIRTestTransforms
   DEPENDS
   MLIRStandardOpsIncGen
 
   LINK_LIBS PUBLIC
   MLIRAffine
   MLIRAnalysis
   MLIREDSC
   MLIRGPU
+  MLIRGPU
   MLIRGPUToGPURuntimeTransforms
   MLIRLinalg
   MLIRLinalgTransforms
   MLIRMathTransforms
   MLIRNVVMIR
-  MLIRSCF
-  MLIRSCFTransforms
-  MLIRGPU
+  MLIRNVVMToLLVMIRTranslation
   MLIRPass
   MLIRROCDLIR
+  MLIRROCDLToLLVMIRTranslation
+  MLIRSCF
+  MLIRSCFTransforms
   MLIRStandardOpsTransforms
-  MLIRTargetNVVMIR
-  MLIRTargetROCDLIR
+  MLIRTargetLLVMIR
   MLIRTestDialect
   MLIRTransformUtils
-  MLIRVectorToSCF
   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/NVVMIR.h"
+#include "mlir/Target/LLVMIR.h"
+#include "mlir/Target/LLVMIR/Dialect/NVVM/NVVMToLLVMIRTranslation.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);
+}
+
 namespace mlir {
 namespace test {
 void registerTestConvertGPUKernelToCubinPass() {
   PassPipelineRegistration<>(
       "test-kernel-to-cubin",
       "Convert all kernel functions to CUDA cubin blobs",
       [](OpPassManager &pm) {
         // Initialize LLVM NVPTX backend.

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/ROCDLIR.h"
+#include "mlir/Target/LLVMIR.h"
+#include "mlir/Target/LLVMIR/Dialect/ROCDL/ROCDLToLLVMIRTranslation.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);
+}
+
 namespace mlir {
 namespace test {
 void registerTestConvertGPUKernelToHsacoPass() {
   PassPipelineRegistration<>(
       "test-kernel-to-hsaco",
       "Convert all kernel functions to ROCm hsaco blobs",
       [](OpPassManager &pm) {
         // Initialize LLVM AMDGPU backend.
         LLVMInitializeAMDGPUTarget();
         LLVMInitializeAMDGPUTargetInfo();
         LLVMInitializeAMDGPUTargetMC();
         LLVMInitializeAMDGPUAsmPrinter();
 
         pm.addPass(createConvertGPUKernelToBlobPass(
-            translateModuleToROCDLIR, compileIsaToHsacoForTesting,
+            translateModuleToROCDL, compileIsaToHsacoForTesting,
             "amdgcn-amd-amdhsa", "gfx900", "-code-object-v3", "rocdl.hsaco"));
       });
 }
 } // namespace test
 } // namespace mlir
 #endif

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

 #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/NVVMIR.h"
+#include "mlir/Target/LLVMIR/Dialect/NVVM/NVVMToLLVMIRTranslation.h"
 #include "mlir/Transforms/DialectConversion.h"
 #include "mlir/Transforms/Passes.h"
 
 #include "llvm/Support/InitLLVM.h"
 #include "llvm/Support/TargetSelect.h"
 
 #include "cuda.h"
 
static LogicalResult runMLIRPasses(ModuleOp m) {
   applyPassManagerCLOptions(pm);
 
   const char gpuBinaryAnnotation[] = "nvvm.cubin";
   pm.addPass(createGpuKernelOutliningPass());
   auto &kernelPm = pm.nest<gpu::GPUModuleOp>();
   kernelPm.addPass(createStripDebugInfoPass());
   kernelPm.addPass(createLowerGpuOpsToNVVMOpsPass());
   kernelPm.addPass(createConvertGPUKernelToBlobPass(
-      translateModuleToNVVMIR, compilePtxToCubin, "nvptx64-nvidia-cuda",
+      translateModuleToLLVMIR, compilePtxToCubin, "nvptx64-nvidia-cuda",
       "sm_35", "+ptx60", gpuBinaryAnnotation));
   auto &funcPm = pm.nest<FuncOp>();
   funcPm.addPass(createGpuAsyncRegionPass());
   funcPm.addPass(createAsyncRefCountingPass());
   pm.addPass(createGpuToLLVMConversionPass(gpuBinaryAnnotation));
   pm.addPass(createAsyncToAsyncRuntimePass());
   pm.addPass(createConvertAsyncToLLVMPass());
   mlir::LowerToLLVMOptions lower_to_llvm_opts;
int main(int argc, char **argv) {
 
   mlir::JitRunnerConfig jitRunnerConfig;
   jitRunnerConfig.mlirTransformer = runMLIRPasses;
 
   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);
 
   return mlir::JitRunnerMain(argc, argv, registry, jitRunnerConfig);
 }

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

 #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/ROCDL/ROCDLToLLVMIRTranslation.h"
 #include "mlir/Target/ROCDLIR.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"
 #include "llvm/Support/Program.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);
 
   return mlir::JitRunnerMain(argc, argv, registry, jitRunnerConfig);
 }

mlir/tools/mlir-tblgen/LLVMIRConversionGen.cpp

static bool emitOneBuilder(const Record &record, raw_ostream &os) {
   llvm::raw_string_ostream bs(builder);
   while (auto loc = findNextVariable(builderStrRef)) {
     auto name = loc.in(builderStrRef).drop_front();
     // First, insert the non-matched part as is.
     bs << builderStrRef.substr(0, loc.pos);
     // Then, rewrite the name based on its kind.
     bool isVariadicOperand = isVariadicOperandName(op, name);
     if (isOperandName(op, name)) {
-      auto result = isVariadicOperand ? formatv("lookupValues(op.{0}())", name)
-                                      : formatv("lookupValue(op.{0}())", name);
+      auto result =
+          isVariadicOperand
+              ? formatv("moduleTranslation.lookupValues(op.{0}())", name)
+              : formatv("moduleTranslation.lookupValue(op.{0}())", name);
       bs << result;
     } else if (isAttributeName(op, name)) {
       bs << formatv("op.{0}()", name);
     } else if (isResultName(op, name)) {
-      bs << formatv("mapValue(op.{0}())", name);
+      bs << formatv("moduleTranslation.mapValue(op.{0}())", name);
     } else if (name == "_resultType") {
-      bs << "convertType(op.getResult().getType())";
+      bs << "moduleTranslation.convertType(op.getResult().getType())";
     } else if (name == "_hasResult") {
       bs << "opInst.getNumResults() == 1";
     } else if (name == "_location") {
       bs << "opInst.getLoc()";
     } else if (name == "_numOperands") {
       bs << "opInst.getNumOperands()";
     } else if (name == "$") {
       bs << '$';