[Mlir-commits] [mlir] 19db802 - [mlir] make implementations of translation to LLVM IR interfaces private
Alex Zinenko
llvmlistbot at llvm.org
Thu Mar 4 00:16:41 PST 2021
Author: Alex Zinenko
Date: 2021-03-04T09:16:32+01:00
New Revision: 19db802e7b13f190ebae6f5389b0bfbc1827e4e7
URL: https://github.com/llvm/llvm-project/commit/19db802e7b13f190ebae6f5389b0bfbc1827e4e7
DIFF: https://github.com/llvm/llvm-project/commit/19db802e7b13f190ebae6f5389b0bfbc1827e4e7.diff
LOG: [mlir] make implementations of translation to LLVM IR interfaces private
There is no need for the interface implementations to be exposed, opaque
registration functions are sufficient for all users, similarly to passes.
Reviewed By: mehdi_amini
Differential Revision: https://reviews.llvm.org/D97852
Added:
mlir/include/mlir/Target/LLVMIR/Dialect/All.h
Modified:
flang/lib/Optimizer/CMakeLists.txt
mlir/examples/toy/Ch6/CMakeLists.txt
mlir/examples/toy/Ch6/toyc.cpp
mlir/examples/toy/Ch7/CMakeLists.txt
mlir/examples/toy/Ch7/toyc.cpp
mlir/include/mlir/Target/LLVMIR.h
mlir/include/mlir/Target/LLVMIR/Dialect/LLVMAVX512/LLVMAVX512ToLLVMIRTranslation.h
mlir/include/mlir/Target/LLVMIR/Dialect/LLVMArmNeon/LLVMArmNeonToLLVMIRTranslation.h
mlir/include/mlir/Target/LLVMIR/Dialect/LLVMArmSVE/LLVMArmSVEToLLVMIRTranslation.h
mlir/include/mlir/Target/LLVMIR/Dialect/LLVMIR/LLVMToLLVMIRTranslation.h
mlir/include/mlir/Target/LLVMIR/Dialect/NVVM/NVVMToLLVMIRTranslation.h
mlir/include/mlir/Target/LLVMIR/Dialect/OpenMP/OpenMPToLLVMIRTranslation.h
mlir/include/mlir/Target/LLVMIR/Dialect/ROCDL/ROCDLToLLVMIRTranslation.h
mlir/include/mlir/Target/LLVMIR/TypeTranslation.h
mlir/lib/CAPI/ExecutionEngine/CMakeLists.txt
mlir/lib/CAPI/ExecutionEngine/ExecutionEngine.cpp
mlir/lib/CAPI/Registration/CMakeLists.txt
mlir/lib/CAPI/Registration/Registration.cpp
mlir/lib/Conversion/VectorToLLVM/CMakeLists.txt
mlir/lib/ExecutionEngine/CMakeLists.txt
mlir/lib/Target/CMakeLists.txt
mlir/lib/Target/LLVMIR/CMakeLists.txt
mlir/lib/Target/LLVMIR/ConvertFromLLVMIR.cpp
mlir/lib/Target/LLVMIR/ConvertToLLVMIR.cpp
mlir/lib/Target/LLVMIR/Dialect/LLVMAVX512/CMakeLists.txt
mlir/lib/Target/LLVMIR/Dialect/LLVMAVX512/LLVMAVX512ToLLVMIRTranslation.cpp
mlir/lib/Target/LLVMIR/Dialect/LLVMArmNeon/CMakeLists.txt
mlir/lib/Target/LLVMIR/Dialect/LLVMArmNeon/LLVMArmNeonToLLVMIRTranslation.cpp
mlir/lib/Target/LLVMIR/Dialect/LLVMArmSVE/CMakeLists.txt
mlir/lib/Target/LLVMIR/Dialect/LLVMArmSVE/LLVMArmSVEToLLVMIRTranslation.cpp
mlir/lib/Target/LLVMIR/Dialect/LLVMIR/CMakeLists.txt
mlir/lib/Target/LLVMIR/Dialect/LLVMIR/LLVMToLLVMIRTranslation.cpp
mlir/lib/Target/LLVMIR/Dialect/NVVM/CMakeLists.txt
mlir/lib/Target/LLVMIR/Dialect/NVVM/NVVMToLLVMIRTranslation.cpp
mlir/lib/Target/LLVMIR/Dialect/OpenMP/CMakeLists.txt
mlir/lib/Target/LLVMIR/Dialect/OpenMP/OpenMPToLLVMIRTranslation.cpp
mlir/lib/Target/LLVMIR/Dialect/ROCDL/CMakeLists.txt
mlir/lib/Target/LLVMIR/Dialect/ROCDL/ROCDLToLLVMIRTranslation.cpp
mlir/lib/Target/LLVMIR/TypeTranslation.cpp
mlir/test/lib/Transforms/CMakeLists.txt
mlir/test/lib/Transforms/TestConvertGPUKernelToCubin.cpp
mlir/test/lib/Transforms/TestConvertGPUKernelToHsaco.cpp
mlir/tools/mlir-cpu-runner/CMakeLists.txt
mlir/tools/mlir-cpu-runner/mlir-cpu-runner.cpp
mlir/tools/mlir-cuda-runner/CMakeLists.txt
mlir/tools/mlir-cuda-runner/mlir-cuda-runner.cpp
mlir/tools/mlir-rocm-runner/CMakeLists.txt
mlir/tools/mlir-rocm-runner/mlir-rocm-runner.cpp
mlir/tools/mlir-spirv-cpu-runner/CMakeLists.txt
mlir/tools/mlir-spirv-cpu-runner/mlir-spirv-cpu-runner.cpp
mlir/tools/mlir-vulkan-runner/CMakeLists.txt
mlir/tools/mlir-vulkan-runner/mlir-vulkan-runner.cpp
mlir/unittests/ExecutionEngine/Invoke.cpp
Removed:
################################################################################
diff --git a/flang/lib/Optimizer/CMakeLists.txt b/flang/lib/Optimizer/CMakeLists.txt
index 9088d08f16c9..0a7286339e2e 100644
--- a/flang/lib/Optimizer/CMakeLists.txt
+++ b/flang/lib/Optimizer/CMakeLists.txt
@@ -19,8 +19,8 @@ add_flang_library(FIROptimizer
LINK_LIBS
${dialect_libs}
- MLIRTargetLLVMIR
- MLIRTargetLLVMIRModuleTranslation
+ MLIRLLVMToLLVMIRTranslation
+ MLIRTargetLLVMIRExport
LINK_COMPONENTS
AsmParser
diff --git a/mlir/examples/toy/Ch6/CMakeLists.txt b/mlir/examples/toy/Ch6/CMakeLists.txt
index 99041c13f22b..62579658e9af 100644
--- a/mlir/examples/toy/Ch6/CMakeLists.txt
+++ b/mlir/examples/toy/Ch6/CMakeLists.txt
@@ -43,10 +43,11 @@ target_link_libraries(toyc-ch6
MLIRExecutionEngine
MLIRIR
MLIRLLVMIR
+ MLIRLLVMToLLVMIRTranslation
MLIRParser
MLIRPass
MLIRSideEffectInterfaces
MLIRSupport
- MLIRTargetLLVMIR
+ MLIRTargetLLVMIRExport
MLIRTransforms
)
diff --git a/mlir/examples/toy/Ch6/toyc.cpp b/mlir/examples/toy/Ch6/toyc.cpp
index 0645bd40c1ba..742a2014c2c0 100644
--- a/mlir/examples/toy/Ch6/toyc.cpp
+++ b/mlir/examples/toy/Ch6/toyc.cpp
@@ -25,7 +25,7 @@
#include "mlir/Parser.h"
#include "mlir/Pass/Pass.h"
#include "mlir/Pass/PassManager.h"
-#include "mlir/Target/LLVMIR.h"
+#include "mlir/Target/LLVMIR/Dialect/LLVMIR/LLVMToLLVMIRTranslation.h"
#include "mlir/Target/LLVMIR/Export.h"
#include "mlir/Transforms/Passes.h"
diff --git a/mlir/examples/toy/Ch7/CMakeLists.txt b/mlir/examples/toy/Ch7/CMakeLists.txt
index c24ad53144e1..cda24c48a5aa 100644
--- a/mlir/examples/toy/Ch7/CMakeLists.txt
+++ b/mlir/examples/toy/Ch7/CMakeLists.txt
@@ -42,9 +42,10 @@ target_link_libraries(toyc-ch7
MLIRCastInterfaces
MLIRExecutionEngine
MLIRIR
+ MLIRLLVMToLLVMIRTranslation
MLIRParser
MLIRPass
MLIRSideEffectInterfaces
- MLIRTargetLLVMIR
+ MLIRTargetLLVMIRExport
MLIRTransforms
)
diff --git a/mlir/examples/toy/Ch7/toyc.cpp b/mlir/examples/toy/Ch7/toyc.cpp
index 0f6d3876a5ad..7218ea0d7277 100644
--- a/mlir/examples/toy/Ch7/toyc.cpp
+++ b/mlir/examples/toy/Ch7/toyc.cpp
@@ -25,7 +25,7 @@
#include "mlir/Parser.h"
#include "mlir/Pass/Pass.h"
#include "mlir/Pass/PassManager.h"
-#include "mlir/Target/LLVMIR.h"
+#include "mlir/Target/LLVMIR/Dialect/LLVMIR/LLVMToLLVMIRTranslation.h"
#include "mlir/Target/LLVMIR/Export.h"
#include "mlir/Transforms/Passes.h"
diff --git a/mlir/include/mlir/Target/LLVMIR.h b/mlir/include/mlir/Target/LLVMIR.h
index 385f55816697..0e801c866837 100644
--- a/mlir/include/mlir/Target/LLVMIR.h
+++ b/mlir/include/mlir/Target/LLVMIR.h
@@ -36,15 +36,6 @@ OwningModuleRef
translateLLVMIRToModule(std::unique_ptr<llvm::Module> llvmModule,
MLIRContext *context);
-/// Register the LLVM dialect and the translation from it to the LLVM IR in the
-/// given registry;
-void registerLLVMDialectTranslation(DialectRegistry ®istry);
-
-/// Register the LLVM dialect and the translation from it in the registry
-/// associated with the given context. This checks if the interface is already
-/// registered and avoids double registation.
-void registerLLVMDialectTranslation(MLIRContext &context);
-
} // namespace mlir
#endif // MLIR_TARGET_LLVMIR_H
diff --git a/mlir/include/mlir/Target/LLVMIR/Dialect/All.h b/mlir/include/mlir/Target/LLVMIR/Dialect/All.h
new file mode 100644
index 000000000000..21df7c85fcab
--- /dev/null
+++ b/mlir/include/mlir/Target/LLVMIR/Dialect/All.h
@@ -0,0 +1,41 @@
+//===- All.h - MLIR To LLVM IR Translation Registration ---------*- C++ -*-===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+//
+// This file defines a helper to register the translations of all suitable
+// dialects to LLVM IR.
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef MLIR_TARGET_LLVMIR_DIALECT_ALL_H
+#define MLIR_TARGET_LLVMIR_DIALECT_ALL_H
+
+#include "mlir/Target/LLVMIR/Dialect/LLVMAVX512/LLVMAVX512ToLLVMIRTranslation.h"
+#include "mlir/Target/LLVMIR/Dialect/LLVMArmNeon/LLVMArmNeonToLLVMIRTranslation.h"
+#include "mlir/Target/LLVMIR/Dialect/LLVMArmSVE/LLVMArmSVEToLLVMIRTranslation.h"
+#include "mlir/Target/LLVMIR/Dialect/LLVMIR/LLVMToLLVMIRTranslation.h"
+#include "mlir/Target/LLVMIR/Dialect/NVVM/NVVMToLLVMIRTranslation.h"
+#include "mlir/Target/LLVMIR/Dialect/OpenMP/OpenMPToLLVMIRTranslation.h"
+#include "mlir/Target/LLVMIR/Dialect/ROCDL/ROCDLToLLVMIRTranslation.h"
+
+namespace mlir {
+class DialectRegistry;
+
+/// Registers all dialects that can be translated to LLVM IR and the
+/// corresponding translation interfaces.
+static inline void registerAllToLLVMIRTranslations(DialectRegistry ®istry) {
+ registerLLVMAVX512DialectTranslation(registry);
+ registerLLVMArmNeonDialectTranslation(registry);
+ registerLLVMArmSVEDialectTranslation(registry);
+ registerLLVMDialectTranslation(registry);
+ registerNVVMDialectTranslation(registry);
+ registerOpenMPDialectTranslation(registry);
+ registerROCDLDialectTranslation(registry);
+}
+} // namespace mlir
+
+#endif // MLIR_TARGET_LLVMIR_DIALECT_ALL_H
diff --git a/mlir/include/mlir/Target/LLVMIR/Dialect/LLVMAVX512/LLVMAVX512ToLLVMIRTranslation.h b/mlir/include/mlir/Target/LLVMIR/Dialect/LLVMAVX512/LLVMAVX512ToLLVMIRTranslation.h
index e591a95f0357..7a871c66edf0 100644
--- a/mlir/include/mlir/Target/LLVMIR/Dialect/LLVMAVX512/LLVMAVX512ToLLVMIRTranslation.h
+++ b/mlir/include/mlir/Target/LLVMIR/Dialect/LLVMAVX512/LLVMAVX512ToLLVMIRTranslation.h
@@ -6,31 +6,26 @@
//
//===----------------------------------------------------------------------===//
//
-// This file implements the dialect interface for translating the LLVMAVX512
-// dialect to LLVM IR.
+// This provides registration calls for LLVMAVX512 dialect to LLVM IR
+// translation.
//
//===----------------------------------------------------------------------===//
#ifndef MLIR_TARGET_LLVMIR_DIALECT_LLVMAVX512_LLVMAVX512TOLLVMIRTRANSLATION_H
#define MLIR_TARGET_LLVMIR_DIALECT_LLVMAVX512_LLVMAVX512TOLLVMIRTRANSLATION_H
-#include "mlir/Target/LLVMIR/LLVMTranslationInterface.h"
-
namespace mlir {
-/// Implementation of the dialect interface that converts operations belonging
-/// to the LLVMAVX512 dialect to LLVM IR.
-class LLVMAVX512DialectLLVMIRTranslationInterface
- : public LLVMTranslationDialectInterface {
-public:
- using LLVMTranslationDialectInterface::LLVMTranslationDialectInterface;
+class DialectRegistry;
+class MLIRContext;
+
+/// Register the LLVMAVX512 dialect and the translation from it to the LLVM IR
+/// in the given registry;
+void registerLLVMAVX512DialectTranslation(DialectRegistry ®istry);
- /// 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;
-};
+/// Register the LLVMAVX512 dialect and the translation from it in the registry
+/// associated with the given context.
+void registerLLVMAVX512DialectTranslation(MLIRContext &context);
} // namespace mlir
diff --git a/mlir/include/mlir/Target/LLVMIR/Dialect/LLVMArmNeon/LLVMArmNeonToLLVMIRTranslation.h b/mlir/include/mlir/Target/LLVMIR/Dialect/LLVMArmNeon/LLVMArmNeonToLLVMIRTranslation.h
index 7d268d155083..daa7441e9711 100644
--- a/mlir/include/mlir/Target/LLVMIR/Dialect/LLVMArmNeon/LLVMArmNeonToLLVMIRTranslation.h
+++ b/mlir/include/mlir/Target/LLVMIR/Dialect/LLVMArmNeon/LLVMArmNeonToLLVMIRTranslation.h
@@ -6,31 +6,26 @@
//
//===----------------------------------------------------------------------===//
//
-// This file implements the dialect interface for translating the LLVMArmNeon
-// dialect to LLVM IR.
+// This provides registration calls for LLVMArmNeon dialect to LLVM IR
+// translation.
//
//===----------------------------------------------------------------------===//
#ifndef MLIR_TARGET_LLVMIR_DIALECT_LLVMARMNEON_LLVMARMNEONTOLLVMIRTRANSLATION_H
#define MLIR_TARGET_LLVMIR_DIALECT_LLVMARMNEON_LLVMARMNEONTOLLVMIRTRANSLATION_H
-#include "mlir/Target/LLVMIR/LLVMTranslationInterface.h"
-
namespace mlir {
-/// Implementation of the dialect interface that converts operations belonging
-/// to the LLVMArmNeon dialect to LLVM IR.
-class LLVMArmNeonDialectLLVMIRTranslationInterface
- : public LLVMTranslationDialectInterface {
-public:
- using LLVMTranslationDialectInterface::LLVMTranslationDialectInterface;
+class DialectRegistry;
+class MLIRContext;
+
+/// Register the LLVMArmNeon dialect and the translation from it to the LLVM IR
+/// in the given registry;
+void registerLLVMArmNeonDialectTranslation(DialectRegistry ®istry);
- /// 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;
-};
+/// Register the LLVMArmNeon dialect and the translation from it in the registry
+/// associated with the given context.
+void registerLLVMArmNeonDialectTranslation(MLIRContext &context);
} // namespace mlir
diff --git a/mlir/include/mlir/Target/LLVMIR/Dialect/LLVMArmSVE/LLVMArmSVEToLLVMIRTranslation.h b/mlir/include/mlir/Target/LLVMIR/Dialect/LLVMArmSVE/LLVMArmSVEToLLVMIRTranslation.h
index 9d4d05b9b9bd..aae865d63cff 100644
--- a/mlir/include/mlir/Target/LLVMIR/Dialect/LLVMArmSVE/LLVMArmSVEToLLVMIRTranslation.h
+++ b/mlir/include/mlir/Target/LLVMIR/Dialect/LLVMArmSVE/LLVMArmSVEToLLVMIRTranslation.h
@@ -6,31 +6,26 @@
//
//===----------------------------------------------------------------------===//
//
-// This file implements the dialect interface for translating the LLVMArmSVE
-// dialect to LLVM IR.
+// This provides registration calls for LLVMArmSVE dialect to LLVM IR
+// translation.
//
//===----------------------------------------------------------------------===//
#ifndef MLIR_TARGET_LLVMIR_DIALECT_LLVMARMSVE_LLVMARMSVETOLLVMIRTRANSLATION_H
#define MLIR_TARGET_LLVMIR_DIALECT_LLVMARMSVE_LLVMARMSVETOLLVMIRTRANSLATION_H
-#include "mlir/Target/LLVMIR/LLVMTranslationInterface.h"
-
namespace mlir {
-/// Implementation of the dialect interface that converts operations belonging
-/// to the LLVMArmSVE dialect to LLVM IR.
-class LLVMArmSVEDialectLLVMIRTranslationInterface
- : public LLVMTranslationDialectInterface {
-public:
- using LLVMTranslationDialectInterface::LLVMTranslationDialectInterface;
+class DialectRegistry;
+class MLIRContext;
+
+/// Register the LLVMArmSVE dialect and the translation from it to the LLVM IR
+/// in the given registry;
+void registerLLVMArmSVEDialectTranslation(DialectRegistry ®istry);
- /// 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;
-};
+/// Register the LLVMArmSVE dialect and the translation from it in the registry
+/// associated with the given context.
+void registerLLVMArmSVEDialectTranslation(MLIRContext &context);
} // namespace mlir
diff --git a/mlir/include/mlir/Target/LLVMIR/Dialect/LLVMIR/LLVMToLLVMIRTranslation.h b/mlir/include/mlir/Target/LLVMIR/Dialect/LLVMIR/LLVMToLLVMIRTranslation.h
index 2af76e092917..8ecb035ea4aa 100644
--- a/mlir/include/mlir/Target/LLVMIR/Dialect/LLVMIR/LLVMToLLVMIRTranslation.h
+++ b/mlir/include/mlir/Target/LLVMIR/Dialect/LLVMIR/LLVMToLLVMIRTranslation.h
@@ -6,31 +6,25 @@
//
//===----------------------------------------------------------------------===//
//
-// This file implements the dialect interface for translating the LLVM dialect
-// to LLVM IR.
+// This provides registration calls for LLVM dialect to LLVM IR translation.
//
//===----------------------------------------------------------------------===//
#ifndef MLIR_TARGET_LLVMIR_DIALECT_LLVMIR_LLVMTOLLVMIRTRANSLATION_H
#define MLIR_TARGET_LLVMIR_DIALECT_LLVMIR_LLVMTOLLVMIRTRANSLATION_H
-#include "mlir/Target/LLVMIR/LLVMTranslationInterface.h"
-
namespace mlir {
-/// Implementation of the dialect interface that converts operations belonging
-/// to the LLVM dialect to LLVM IR.
-class LLVMDialectLLVMIRTranslationInterface
- : public LLVMTranslationDialectInterface {
-public:
- using LLVMTranslationDialectInterface::LLVMTranslationDialectInterface;
+class DialectRegistry;
+class MLIRContext;
+
+/// Register the LLVM dialect and the translation from it to the LLVM IR in the
+/// given registry;
+void registerLLVMDialectTranslation(DialectRegistry ®istry);
- /// 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;
-};
+/// Register the LLVM dialect and the translation from it in the registry
+/// associated with the given context.
+void registerLLVMDialectTranslation(MLIRContext &context);
} // namespace mlir
diff --git a/mlir/include/mlir/Target/LLVMIR/Dialect/NVVM/NVVMToLLVMIRTranslation.h b/mlir/include/mlir/Target/LLVMIR/Dialect/NVVM/NVVMToLLVMIRTranslation.h
index 3a8a01df84b0..723bba678b08 100644
--- a/mlir/include/mlir/Target/LLVMIR/Dialect/NVVM/NVVMToLLVMIRTranslation.h
+++ b/mlir/include/mlir/Target/LLVMIR/Dialect/NVVM/NVVMToLLVMIRTranslation.h
@@ -6,36 +6,25 @@
//
//===----------------------------------------------------------------------===//
//
-// This file implements the dialect interface for translating the NVVM
-// dialect to LLVM IR.
+// This provides registration calls for NVVM dialect to LLVM IR translation.
//
//===----------------------------------------------------------------------===//
#ifndef MLIR_TARGET_LLVMIR_DIALECT_NVVM_NVVMTOLLVMIRTRANSLATION_H
#define MLIR_TARGET_LLVMIR_DIALECT_NVVM_NVVMTOLLVMIRTRANSLATION_H
-#include "mlir/Target/LLVMIR/LLVMTranslationInterface.h"
-
namespace mlir {
-/// Implementation of the dialect interface that converts operations belonging
-/// to the NVVM dialect to LLVM IR.
-class NVVMDialectLLVMIRTranslationInterface
- : public LLVMTranslationDialectInterface {
-public:
- using LLVMTranslationDialectInterface::LLVMTranslationDialectInterface;
+class DialectRegistry;
+class MLIRContext;
- /// 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;
+/// Register the NVVM dialect and the translation from it to the LLVM IR in the
+/// given registry;
+void registerNVVMDialectTranslation(DialectRegistry ®istry);
- /// Attaches module-level metadata for functions marked as kernels.
- LogicalResult
- amendOperation(Operation *op, NamedAttribute attribute,
- LLVM::ModuleTranslation &moduleTranslation) const final;
-};
+/// Register the NVVM dialect and the translation from it in the registry
+/// associated with the given context.
+void registerNVVMDialectTranslation(MLIRContext &context);
} // namespace mlir
diff --git a/mlir/include/mlir/Target/LLVMIR/Dialect/OpenMP/OpenMPToLLVMIRTranslation.h b/mlir/include/mlir/Target/LLVMIR/Dialect/OpenMP/OpenMPToLLVMIRTranslation.h
index 07721d089689..bda81227fa23 100644
--- a/mlir/include/mlir/Target/LLVMIR/Dialect/OpenMP/OpenMPToLLVMIRTranslation.h
+++ b/mlir/include/mlir/Target/LLVMIR/Dialect/OpenMP/OpenMPToLLVMIRTranslation.h
@@ -6,31 +6,25 @@
//
//===----------------------------------------------------------------------===//
//
-// This file implements the dialect interface for translating the OpenMP dialect
-// to LLVM IR.
+// This provides registration calls for OpenMP dialect to LLVM IR translation.
//
//===----------------------------------------------------------------------===//
#ifndef MLIR_TARGET_LLVMIR_DIALECT_OPENMP_OPENMPTOLLVMIRTRANSLATION_H
#define MLIR_TARGET_LLVMIR_DIALECT_OPENMP_OPENMPTOLLVMIRTRANSLATION_H
-#include "mlir/Target/LLVMIR/LLVMTranslationInterface.h"
-
namespace mlir {
-/// Implementation of the dialect interface that converts operations belonging
-/// to the OpenMP dialect to LLVM IR.
-class OpenMPDialectLLVMIRTranslationInterface
- : public LLVMTranslationDialectInterface {
-public:
- using LLVMTranslationDialectInterface::LLVMTranslationDialectInterface;
+class DialectRegistry;
+class MLIRContext;
+
+/// Register the OpenMP dialect and the translation from it to the LLVM IR in
+/// the given registry;
+void registerOpenMPDialectTranslation(DialectRegistry ®istry);
- /// 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;
-};
+/// Register the OpenMP dialect and the translation from it in the registry
+/// associated with the given context.
+void registerOpenMPDialectTranslation(MLIRContext &context);
} // namespace mlir
diff --git a/mlir/include/mlir/Target/LLVMIR/Dialect/ROCDL/ROCDLToLLVMIRTranslation.h b/mlir/include/mlir/Target/LLVMIR/Dialect/ROCDL/ROCDLToLLVMIRTranslation.h
index e2211a59098f..ad9061e8f725 100644
--- a/mlir/include/mlir/Target/LLVMIR/Dialect/ROCDL/ROCDLToLLVMIRTranslation.h
+++ b/mlir/include/mlir/Target/LLVMIR/Dialect/ROCDL/ROCDLToLLVMIRTranslation.h
@@ -6,36 +6,25 @@
//
//===----------------------------------------------------------------------===//
//
-// This file implements the dialect interface for translating the ROCDL
-// dialect to LLVM IR.
+// This provides registration calls for ROCDL dialect to LLVM IR translation.
//
//===----------------------------------------------------------------------===//
#ifndef MLIR_TARGET_LLVMIR_DIALECT_ROCDL_ROCDLTOLLVMIRTRANSLATION_H
#define MLIR_TARGET_LLVMIR_DIALECT_ROCDL_ROCDLTOLLVMIRTRANSLATION_H
-#include "mlir/Target/LLVMIR/LLVMTranslationInterface.h"
-
namespace mlir {
-/// Implementation of the dialect interface that converts operations belonging
-/// to the ROCDL dialect to LLVM IR.
-class ROCDLDialectLLVMIRTranslationInterface
- : public LLVMTranslationDialectInterface {
-public:
- using LLVMTranslationDialectInterface::LLVMTranslationDialectInterface;
+class DialectRegistry;
+class MLIRContext;
- /// 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;
+/// Register the ROCDL dialect and the translation from it to the LLVM IR in the
+/// given registry;
+void registerROCDLDialectTranslation(DialectRegistry ®istry);
- /// Attaches module-level metadata for functions marked as kernels.
- LogicalResult
- amendOperation(Operation *op, NamedAttribute attribute,
- LLVM::ModuleTranslation &moduleTranslation) const final;
-};
+/// Register the ROCDL dialect and the translation from it in the registry
+/// associated with the given context.
+void registerROCDLDialectTranslation(MLIRContext &context);
} // namespace mlir
diff --git a/mlir/include/mlir/Target/LLVMIR/TypeTranslation.h b/mlir/include/mlir/Target/LLVMIR/TypeTranslation.h
index c030d51305e4..893167e33105 100644
--- a/mlir/include/mlir/Target/LLVMIR/TypeTranslation.h
+++ b/mlir/include/mlir/Target/LLVMIR/TypeTranslation.h
@@ -31,7 +31,6 @@ namespace LLVM {
namespace detail {
class TypeToLLVMIRTranslatorImpl;
-class TypeFromLLVMIRTranslatorImpl;
} // namespace detail
/// Utility class to translate MLIR LLVM dialect types to LLVM IR. Stores the
@@ -56,22 +55,6 @@ class TypeToLLVMIRTranslator {
std::unique_ptr<detail::TypeToLLVMIRTranslatorImpl> impl;
};
-/// Utility class to translate LLVM IR types to the MLIR LLVM dialect. Stores
-/// the translation state, in particular any identified structure types that are
-/// reused across translations.
-class TypeFromLLVMIRTranslator {
-public:
- TypeFromLLVMIRTranslator(MLIRContext &context);
- ~TypeFromLLVMIRTranslator();
-
- /// Translates the given LLVM IR type to the MLIR LLVM dialect.
- Type translateType(llvm::Type *type);
-
-private:
- /// Private implementation.
- std::unique_ptr<detail::TypeFromLLVMIRTranslatorImpl> impl;
-};
-
} // namespace LLVM
} // namespace mlir
diff --git a/mlir/lib/CAPI/ExecutionEngine/CMakeLists.txt b/mlir/lib/CAPI/ExecutionEngine/CMakeLists.txt
index ab1c8e9ddaa9..09dcb6143745 100644
--- a/mlir/lib/CAPI/ExecutionEngine/CMakeLists.txt
+++ b/mlir/lib/CAPI/ExecutionEngine/CMakeLists.txt
@@ -4,4 +4,5 @@ add_mlir_public_c_api_library(MLIRCEXECUTIONENGINE
LINK_LIBS PUBLIC
MLIRExecutionEngine
+ MLIRLLVMToLLVMIRTranslation
)
diff --git a/mlir/lib/CAPI/ExecutionEngine/ExecutionEngine.cpp b/mlir/lib/CAPI/ExecutionEngine/ExecutionEngine.cpp
index 49722e252ad5..68137c0674c5 100644
--- a/mlir/lib/CAPI/ExecutionEngine/ExecutionEngine.cpp
+++ b/mlir/lib/CAPI/ExecutionEngine/ExecutionEngine.cpp
@@ -10,7 +10,7 @@
#include "mlir/CAPI/ExecutionEngine.h"
#include "mlir/CAPI/IR.h"
#include "mlir/CAPI/Support.h"
-#include "mlir/Target/LLVMIR.h"
+#include "mlir/Target/LLVMIR/Dialect/LLVMIR/LLVMToLLVMIRTranslation.h"
#include "llvm/Support/TargetSelect.h"
using namespace mlir;
diff --git a/mlir/lib/CAPI/Registration/CMakeLists.txt b/mlir/lib/CAPI/Registration/CMakeLists.txt
index 3ac29e2a4128..417140ac0ac6 100644
--- a/mlir/lib/CAPI/Registration/CMakeLists.txt
+++ b/mlir/lib/CAPI/Registration/CMakeLists.txt
@@ -5,5 +5,6 @@ add_mlir_public_c_api_library(MLIRCAPIRegistration
LINK_LIBS PUBLIC
MLIRCAPIIR
+ MLIRLLVMToLLVMIRTranslation
${dialect_libs}
)
diff --git a/mlir/lib/CAPI/Registration/Registration.cpp b/mlir/lib/CAPI/Registration/Registration.cpp
index f8e1bd3f3cb1..dea782453152 100644
--- a/mlir/lib/CAPI/Registration/Registration.cpp
+++ b/mlir/lib/CAPI/Registration/Registration.cpp
@@ -10,7 +10,7 @@
#include "mlir/CAPI/IR.h"
#include "mlir/InitAllDialects.h"
-#include "mlir/Target/LLVMIR.h"
+#include "mlir/Target/LLVMIR/Dialect/LLVMIR/LLVMToLLVMIRTranslation.h"
void mlirRegisterAllDialects(MlirContext context) {
mlir::registerAllDialects(*unwrap(context));
diff --git a/mlir/lib/Conversion/VectorToLLVM/CMakeLists.txt b/mlir/lib/Conversion/VectorToLLVM/CMakeLists.txt
index f1d662b630b1..e2123c3287b2 100644
--- a/mlir/lib/Conversion/VectorToLLVM/CMakeLists.txt
+++ b/mlir/lib/Conversion/VectorToLLVM/CMakeLists.txt
@@ -24,7 +24,7 @@ add_mlir_conversion_library(MLIRVectorToLLVM
MLIRLLVMArmSVE
MLIRLLVMIR
MLIRStandardToLLVM
- MLIRTargetLLVMIRModuleTranslation
+ MLIRTargetLLVMIRExport
MLIRTransforms
MLIRVector
)
diff --git a/mlir/lib/ExecutionEngine/CMakeLists.txt b/mlir/lib/ExecutionEngine/CMakeLists.txt
index 1ec6a8cf9698..b9176cf1e89b 100644
--- a/mlir/lib/ExecutionEngine/CMakeLists.txt
+++ b/mlir/lib/ExecutionEngine/CMakeLists.txt
@@ -44,7 +44,8 @@ add_mlir_library(MLIRExecutionEngine
LINK_LIBS PUBLIC
MLIRLLVMIR
- MLIRTargetLLVMIR
+ MLIRLLVMToLLVMIRTranslation
+ MLIRTargetLLVMIRExport
)
get_property(dialect_libs GLOBAL PROPERTY MLIR_DIALECT_LIBS)
@@ -67,7 +68,8 @@ add_mlir_library(MLIRJitRunner
MLIRIR
MLIRParser
MLIRStandard
- MLIRTargetLLVMIR
+ MLIRLLVMToLLVMIRTranslation
+ MLIRTargetLLVMIRExport
MLIRTransforms
MLIRStandardToLLVM
MLIRSupport
diff --git a/mlir/lib/Target/CMakeLists.txt b/mlir/lib/Target/CMakeLists.txt
index a23222d37ede..5c0a142d4cb4 100644
--- a/mlir/lib/Target/CMakeLists.txt
+++ b/mlir/lib/Target/CMakeLists.txt
@@ -1,47 +1,2 @@
add_subdirectory(SPIRV)
add_subdirectory(LLVMIR)
-
-add_mlir_translation_library(MLIRTargetLLVMIRModuleTranslation
- LLVMIR/DebugTranslation.cpp
- LLVMIR/ModuleTranslation.cpp
- LLVMIR/TypeTranslation.cpp
-
- ADDITIONAL_HEADER_DIRS
- ${MLIR_MAIN_INCLUDE_DIR}/mlir/Target/LLVMIR
-
- DEPENDS
- intrinsics_gen
-
- LINK_COMPONENTS
- Core
- FrontendOpenMP
- TransformUtils
-
- LINK_LIBS PUBLIC
- MLIRLLVMIR
- MLIROpenMP
- MLIRLLVMIRTransforms
- MLIRTranslation
- )
-
-add_mlir_translation_library(MLIRTargetLLVMIR
- LLVMIR/ConvertFromLLVMIR.cpp
- LLVMIR/ConvertToLLVMIR.cpp
-
- ADDITIONAL_HEADER_DIRS
- ${MLIR_MAIN_INCLUDE_DIR}/mlir/Target/LLVMIR
-
- LINK_COMPONENTS
- Core
- IRReader
-
- LINK_LIBS PUBLIC
- MLIRLLVMArmNeonToLLVMIRTranslation
- MLIRLLVMArmSVEToLLVMIRTranslation
- MLIRLLVMAVX512ToLLVMIRTranslation
- MLIRLLVMToLLVMIRTranslation
- MLIRNVVMToLLVMIRTranslation
- MLIROpenMPToLLVMIRTranslation
- MLIRROCDLToLLVMIRTranslation
- MLIRTargetLLVMIRModuleTranslation
- )
diff --git a/mlir/lib/Target/LLVMIR/CMakeLists.txt b/mlir/lib/Target/LLVMIR/CMakeLists.txt
index 0ca0f41c5af4..4c6863259639 100644
--- a/mlir/lib/Target/LLVMIR/CMakeLists.txt
+++ b/mlir/lib/Target/LLVMIR/CMakeLists.txt
@@ -1 +1,61 @@
add_subdirectory(Dialect)
+
+set(LLVM_OPTIONAL_SOURCES
+ ConvertFromLLVMIR.cpp
+ ConvertToLLVMIR.cpp
+ DebugTranslation.cpp
+ ModuleTranslation.cpp
+ TypeTranslation.cpp
+ )
+
+
+add_mlir_translation_library(MLIRTargetLLVMIRExport
+ DebugTranslation.cpp
+ ModuleTranslation.cpp
+ TypeTranslation.cpp
+
+ ADDITIONAL_HEADER_DIRS
+ ${MLIR_MAIN_INCLUDE_DIR}/mlir/Target/LLVMIR
+
+ DEPENDS
+ intrinsics_gen
+
+ LINK_COMPONENTS
+ Core
+ FrontendOpenMP
+ TransformUtils
+
+ LINK_LIBS PUBLIC
+ MLIRLLVMIR
+ MLIROpenMP
+ MLIRLLVMIRTransforms
+ MLIRTranslation
+ )
+
+add_mlir_translation_library(MLIRToLLVMIRTranslationRegistration
+ ConvertToLLVMIR.cpp
+
+ LINK_LIBS PUBLIC
+ MLIRLLVMArmNeonToLLVMIRTranslation
+ MLIRLLVMArmSVEToLLVMIRTranslation
+ MLIRLLVMAVX512ToLLVMIRTranslation
+ MLIRLLVMToLLVMIRTranslation
+ MLIRNVVMToLLVMIRTranslation
+ MLIROpenMPToLLVMIRTranslation
+ MLIRROCDLToLLVMIRTranslation
+ )
+
+add_mlir_translation_library(MLIRTargetLLVMIRImport
+ ConvertFromLLVMIR.cpp
+
+ ADDITIONAL_HEADER_DIRS
+ ${MLIR_MAIN_INCLUDE_DIR}/mlir/Target/LLVMIR
+
+ LINK_COMPONENTS
+ Core
+ IRReader
+
+ LINK_LIBS PUBLIC
+ MLIRLLVMIR
+ MLIRTranslation
+ )
diff --git a/mlir/lib/Target/LLVMIR/ConvertFromLLVMIR.cpp b/mlir/lib/Target/LLVMIR/ConvertFromLLVMIR.cpp
index 0a8a11337859..7432451b7a78 100644
--- a/mlir/lib/Target/LLVMIR/ConvertFromLLVMIR.cpp
+++ b/mlir/lib/Target/LLVMIR/ConvertFromLLVMIR.cpp
@@ -16,11 +16,12 @@
#include "mlir/IR/BuiltinTypes.h"
#include "mlir/IR/MLIRContext.h"
#include "mlir/Target/LLVMIR.h"
-#include "mlir/Target/LLVMIR/TypeTranslation.h"
#include "mlir/Translation.h"
+#include "llvm/ADT/TypeSwitch.h"
#include "llvm/IR/Attributes.h"
#include "llvm/IR/Constants.h"
+#include "llvm/IR/DerivedTypes.h"
#include "llvm/IR/Function.h"
#include "llvm/IR/InlineAsm.h"
#include "llvm/IR/Instructions.h"
@@ -44,6 +45,167 @@ static std::string diag(llvm::Value &v) {
return os.str();
}
+namespace mlir {
+namespace LLVM {
+namespace detail {
+/// Support for translating LLVM IR types to MLIR LLVM dialect types.
+class TypeFromLLVMIRTranslatorImpl {
+public:
+ /// Constructs a class creating types in the given MLIR context.
+ TypeFromLLVMIRTranslatorImpl(MLIRContext &context) : context(context) {}
+
+ /// Translates the given type.
+ Type translateType(llvm::Type *type) {
+ if (knownTranslations.count(type))
+ return knownTranslations.lookup(type);
+
+ Type translated =
+ llvm::TypeSwitch<llvm::Type *, Type>(type)
+ .Case<llvm::ArrayType, llvm::FunctionType, llvm::IntegerType,
+ llvm::PointerType, llvm::StructType, llvm::FixedVectorType,
+ llvm::ScalableVectorType>(
+ [this](auto *type) { return this->translate(type); })
+ .Default([this](llvm::Type *type) {
+ return translatePrimitiveType(type);
+ });
+ knownTranslations.try_emplace(type, translated);
+ return translated;
+ }
+
+private:
+ /// Translates the given primitive, i.e. non-parametric in MLIR nomenclature,
+ /// type.
+ Type translatePrimitiveType(llvm::Type *type) {
+ if (type->isVoidTy())
+ return LLVM::LLVMVoidType::get(&context);
+ if (type->isHalfTy())
+ return Float16Type::get(&context);
+ if (type->isBFloatTy())
+ return BFloat16Type::get(&context);
+ if (type->isFloatTy())
+ return Float32Type::get(&context);
+ if (type->isDoubleTy())
+ return Float64Type::get(&context);
+ if (type->isFP128Ty())
+ return Float128Type::get(&context);
+ if (type->isX86_FP80Ty())
+ return Float80Type::get(&context);
+ if (type->isPPC_FP128Ty())
+ return LLVM::LLVMPPCFP128Type::get(&context);
+ if (type->isX86_MMXTy())
+ return LLVM::LLVMX86MMXType::get(&context);
+ if (type->isLabelTy())
+ return LLVM::LLVMLabelType::get(&context);
+ if (type->isMetadataTy())
+ return LLVM::LLVMMetadataType::get(&context);
+ llvm_unreachable("not a primitive type");
+ }
+
+ /// Translates the given array type.
+ Type translate(llvm::ArrayType *type) {
+ return LLVM::LLVMArrayType::get(translateType(type->getElementType()),
+ type->getNumElements());
+ }
+
+ /// Translates the given function type.
+ Type translate(llvm::FunctionType *type) {
+ SmallVector<Type, 8> paramTypes;
+ translateTypes(type->params(), paramTypes);
+ return LLVM::LLVMFunctionType::get(translateType(type->getReturnType()),
+ paramTypes, type->isVarArg());
+ }
+
+ /// Translates the given integer type.
+ Type translate(llvm::IntegerType *type) {
+ return IntegerType::get(&context, type->getBitWidth());
+ }
+
+ /// Translates the given pointer type.
+ Type translate(llvm::PointerType *type) {
+ return LLVM::LLVMPointerType::get(translateType(type->getElementType()),
+ type->getAddressSpace());
+ }
+
+ /// Translates the given structure type.
+ Type translate(llvm::StructType *type) {
+ SmallVector<Type, 8> subtypes;
+ if (type->isLiteral()) {
+ translateTypes(type->subtypes(), subtypes);
+ return LLVM::LLVMStructType::getLiteral(&context, subtypes,
+ type->isPacked());
+ }
+
+ if (type->isOpaque())
+ return LLVM::LLVMStructType::getOpaque(type->getName(), &context);
+
+ LLVM::LLVMStructType translated =
+ LLVM::LLVMStructType::getIdentified(&context, type->getName());
+ knownTranslations.try_emplace(type, translated);
+ translateTypes(type->subtypes(), subtypes);
+ LogicalResult bodySet = translated.setBody(subtypes, type->isPacked());
+ assert(succeeded(bodySet) &&
+ "could not set the body of an identified struct");
+ (void)bodySet;
+ return translated;
+ }
+
+ /// Translates the given fixed-vector type.
+ Type translate(llvm::FixedVectorType *type) {
+ return LLVM::getFixedVectorType(translateType(type->getElementType()),
+ type->getNumElements());
+ }
+
+ /// Translates the given scalable-vector type.
+ Type translate(llvm::ScalableVectorType *type) {
+ return LLVM::LLVMScalableVectorType::get(
+ translateType(type->getElementType()), type->getMinNumElements());
+ }
+
+ /// Translates a list of types.
+ void translateTypes(ArrayRef<llvm::Type *> types,
+ SmallVectorImpl<Type> &result) {
+ result.reserve(result.size() + types.size());
+ for (llvm::Type *type : types)
+ result.push_back(translateType(type));
+ }
+
+ /// Map of known translations. Serves as a cache and as recursion stopper for
+ /// translating recursive structs.
+ llvm::DenseMap<llvm::Type *, Type> knownTranslations;
+
+ /// The context in which MLIR types are created.
+ MLIRContext &context;
+};
+} // end namespace detail
+
+/// Utility class to translate LLVM IR types to the MLIR LLVM dialect. Stores
+/// the translation state, in particular any identified structure types that are
+/// reused across translations.
+class TypeFromLLVMIRTranslator {
+public:
+ TypeFromLLVMIRTranslator(MLIRContext &context);
+ ~TypeFromLLVMIRTranslator();
+
+ /// Translates the given LLVM IR type to the MLIR LLVM dialect.
+ Type translateType(llvm::Type *type);
+
+private:
+ /// Private implementation.
+ std::unique_ptr<detail::TypeFromLLVMIRTranslatorImpl> impl;
+};
+
+} // end namespace LLVM
+} // end namespace mlir
+
+LLVM::TypeFromLLVMIRTranslator::TypeFromLLVMIRTranslator(MLIRContext &context)
+ : impl(new detail::TypeFromLLVMIRTranslatorImpl(context)) {}
+
+LLVM::TypeFromLLVMIRTranslator::~TypeFromLLVMIRTranslator() {}
+
+Type LLVM::TypeFromLLVMIRTranslator::translateType(llvm::Type *type) {
+ return impl->translateType(type);
+}
+
// Handles importing globals and functions from an LLVM module.
namespace {
class Importer {
diff --git a/mlir/lib/Target/LLVMIR/ConvertToLLVMIR.cpp b/mlir/lib/Target/LLVMIR/ConvertToLLVMIR.cpp
index f36b0a41f2ec..ef9de9490956 100644
--- a/mlir/lib/Target/LLVMIR/ConvertToLLVMIR.cpp
+++ b/mlir/lib/Target/LLVMIR/ConvertToLLVMIR.cpp
@@ -10,49 +10,15 @@
//
//===----------------------------------------------------------------------===//
-#include "mlir/Target/LLVMIR.h"
-
-#include "mlir/Dialect/LLVMIR/LLVMAVX512Dialect.h"
-#include "mlir/Dialect/LLVMIR/LLVMArmNeonDialect.h"
-#include "mlir/Dialect/LLVMIR/LLVMArmSVEDialect.h"
-#include "mlir/Dialect/LLVMIR/NVVMDialect.h"
-#include "mlir/Dialect/LLVMIR/ROCDLDialect.h"
-#include "mlir/Dialect/OpenMP/OpenMPDialect.h"
-#include "mlir/Target/LLVMIR/Dialect/LLVMAVX512/LLVMAVX512ToLLVMIRTranslation.h"
-#include "mlir/Target/LLVMIR/Dialect/LLVMArmNeon/LLVMArmNeonToLLVMIRTranslation.h"
-#include "mlir/Target/LLVMIR/Dialect/LLVMArmSVE/LLVMArmSVEToLLVMIRTranslation.h"
-#include "mlir/Target/LLVMIR/Dialect/LLVMIR/LLVMToLLVMIRTranslation.h"
-#include "mlir/Target/LLVMIR/Dialect/NVVM/NVVMToLLVMIRTranslation.h"
-#include "mlir/Target/LLVMIR/Dialect/OpenMP/OpenMPToLLVMIRTranslation.h"
-#include "mlir/Target/LLVMIR/Dialect/ROCDL/ROCDLToLLVMIRTranslation.h"
-#include "mlir/Target/LLVMIR/ModuleTranslation.h"
+#include "mlir/IR/BuiltinOps.h"
+#include "mlir/Target/LLVMIR/Dialect/All.h"
+#include "mlir/Target/LLVMIR/Export.h"
#include "mlir/Translation.h"
-
-#include "llvm/ADT/StringRef.h"
+#include "llvm/IR/LLVMContext.h"
#include "llvm/IR/Module.h"
-#include "llvm/IR/Verifier.h"
-#include "llvm/Support/ToolOutputFile.h"
using namespace mlir;
-void mlir::registerLLVMDialectTranslation(DialectRegistry ®istry) {
- registry.insert<LLVM::LLVMDialect>();
- registry.addDialectInterface<LLVM::LLVMDialect,
- LLVMDialectLLVMIRTranslationInterface>();
-}
-
-void mlir::registerLLVMDialectTranslation(MLIRContext &context) {
- auto *dialect = context.getLoadedDialect<LLVM::LLVMDialect>();
- if (!dialect || dialect->getRegisteredInterface<
- LLVMDialectLLVMIRTranslationInterface>() == nullptr) {
- DialectRegistry registry;
- registry.insert<LLVM::LLVMDialect>();
- registry.addDialectInterface<LLVM::LLVMDialect,
- LLVMDialectLLVMIRTranslationInterface>();
- context.appendDialectRegistry(registry);
- }
-}
-
namespace mlir {
void registerToLLVMIRTranslation() {
TranslateFromMLIRRegistration registration(
@@ -67,25 +33,7 @@ void registerToLLVMIRTranslation() {
return success();
},
[](DialectRegistry ®istry) {
- registry.insert<omp::OpenMPDialect, LLVM::LLVMAVX512Dialect,
- LLVM::LLVMArmSVEDialect, LLVM::LLVMArmNeonDialect,
- NVVM::NVVMDialect, ROCDL::ROCDLDialect>();
- registry.addDialectInterface<omp::OpenMPDialect,
- OpenMPDialectLLVMIRTranslationInterface>();
- registry
- .addDialectInterface<LLVM::LLVMAVX512Dialect,
- LLVMAVX512DialectLLVMIRTranslationInterface>();
- registry.addDialectInterface<
- LLVM::LLVMArmNeonDialect,
- LLVMArmNeonDialectLLVMIRTranslationInterface>();
- registry
- .addDialectInterface<LLVM::LLVMArmSVEDialect,
- LLVMArmSVEDialectLLVMIRTranslationInterface>();
- registry.addDialectInterface<NVVM::NVVMDialect,
- NVVMDialectLLVMIRTranslationInterface>();
- registry.addDialectInterface<ROCDL::ROCDLDialect,
- ROCDLDialectLLVMIRTranslationInterface>();
- registerLLVMDialectTranslation(registry);
+ registerAllToLLVMIRTranslations(registry);
});
}
} // namespace mlir
diff --git a/mlir/lib/Target/LLVMIR/Dialect/LLVMAVX512/CMakeLists.txt b/mlir/lib/Target/LLVMIR/Dialect/LLVMAVX512/CMakeLists.txt
index ba7f594cd81a..3678ab5384ff 100644
--- a/mlir/lib/Target/LLVMIR/Dialect/LLVMAVX512/CMakeLists.txt
+++ b/mlir/lib/Target/LLVMIR/Dialect/LLVMAVX512/CMakeLists.txt
@@ -12,5 +12,5 @@ add_mlir_translation_library(MLIRLLVMAVX512ToLLVMIRTranslation
MLIRLLVMAVX512
MLIRLLVMIR
MLIRSupport
- MLIRTargetLLVMIRModuleTranslation
+ MLIRTargetLLVMIRExport
)
diff --git a/mlir/lib/Target/LLVMIR/Dialect/LLVMAVX512/LLVMAVX512ToLLVMIRTranslation.cpp b/mlir/lib/Target/LLVMIR/Dialect/LLVMAVX512/LLVMAVX512ToLLVMIRTranslation.cpp
index dc4742d8e902..7d1e087dfdb7 100644
--- a/mlir/lib/Target/LLVMIR/Dialect/LLVMAVX512/LLVMAVX512ToLLVMIRTranslation.cpp
+++ b/mlir/lib/Target/LLVMIR/Dialect/LLVMAVX512/LLVMAVX512ToLLVMIRTranslation.cpp
@@ -22,12 +22,35 @@
using namespace mlir;
using namespace mlir::LLVM;
-LogicalResult
-mlir::LLVMAVX512DialectLLVMIRTranslationInterface::convertOperation(
- Operation *op, llvm::IRBuilderBase &builder,
- LLVM::ModuleTranslation &moduleTranslation) const {
- Operation &opInst = *op;
+namespace {
+/// Implementation of the dialect interface that converts operations belonging
+/// to the LLVMAVX512 dialect to LLVM IR.
+class LLVMAVX512DialectLLVMIRTranslationInterface
+ : public LLVMTranslationDialectInterface {
+public:
+ using LLVMTranslationDialectInterface::LLVMTranslationDialectInterface;
+
+ /// Translates the given operation to LLVM IR using the provided IR builder
+ /// and saving the state in `moduleTranslation`.
+ LogicalResult
+ convertOperation(Operation *op, llvm::IRBuilderBase &builder,
+ LLVM::ModuleTranslation &moduleTranslation) const final {
+ Operation &opInst = *op;
#include "mlir/Dialect/LLVMIR/LLVMAVX512Conversions.inc"
- return failure();
+ return failure();
+ }
+};
+} // end namespace
+
+void mlir::registerLLVMAVX512DialectTranslation(DialectRegistry ®istry) {
+ registry.insert<LLVM::LLVMAVX512Dialect>();
+ registry.addDialectInterface<LLVM::LLVMAVX512Dialect,
+ LLVMAVX512DialectLLVMIRTranslationInterface>();
+}
+
+void mlir::registerLLVMAVX512DialectTranslation(MLIRContext &context) {
+ DialectRegistry registry;
+ registerLLVMAVX512DialectTranslation(registry);
+ context.appendDialectRegistry(registry);
}
diff --git a/mlir/lib/Target/LLVMIR/Dialect/LLVMArmNeon/CMakeLists.txt b/mlir/lib/Target/LLVMIR/Dialect/LLVMArmNeon/CMakeLists.txt
index 4ff759ad7521..634e775c4390 100644
--- a/mlir/lib/Target/LLVMIR/Dialect/LLVMArmNeon/CMakeLists.txt
+++ b/mlir/lib/Target/LLVMIR/Dialect/LLVMArmNeon/CMakeLists.txt
@@ -12,5 +12,5 @@ add_mlir_translation_library(MLIRLLVMArmNeonToLLVMIRTranslation
MLIRLLVMArmNeon
MLIRLLVMIR
MLIRSupport
- MLIRTargetLLVMIRModuleTranslation
+ MLIRTargetLLVMIRExport
)
diff --git a/mlir/lib/Target/LLVMIR/Dialect/LLVMArmNeon/LLVMArmNeonToLLVMIRTranslation.cpp b/mlir/lib/Target/LLVMIR/Dialect/LLVMArmNeon/LLVMArmNeonToLLVMIRTranslation.cpp
index 80776ef19cfc..f8a175baa78a 100644
--- a/mlir/lib/Target/LLVMIR/Dialect/LLVMArmNeon/LLVMArmNeonToLLVMIRTranslation.cpp
+++ b/mlir/lib/Target/LLVMIR/Dialect/LLVMArmNeon/LLVMArmNeonToLLVMIRTranslation.cpp
@@ -22,12 +22,35 @@
using namespace mlir;
using namespace mlir::LLVM;
-LogicalResult
-mlir::LLVMArmNeonDialectLLVMIRTranslationInterface::convertOperation(
- Operation *op, llvm::IRBuilderBase &builder,
- LLVM::ModuleTranslation &moduleTranslation) const {
- Operation &opInst = *op;
+namespace {
+/// Implementation of the dialect interface that converts operations belonging
+/// to the LLVMArmNeon dialect to LLVM IR.
+class LLVMArmNeonDialectLLVMIRTranslationInterface
+ : public LLVMTranslationDialectInterface {
+public:
+ using LLVMTranslationDialectInterface::LLVMTranslationDialectInterface;
+
+ /// Translates the given operation to LLVM IR using the provided IR builder
+ /// and saving the state in `moduleTranslation`.
+ LogicalResult
+ convertOperation(Operation *op, llvm::IRBuilderBase &builder,
+ LLVM::ModuleTranslation &moduleTranslation) const final {
+ Operation &opInst = *op;
#include "mlir/Dialect/LLVMIR/LLVMArmNeonConversions.inc"
- return failure();
+ return failure();
+ }
+};
+} // end namespace
+
+void mlir::registerLLVMArmNeonDialectTranslation(DialectRegistry ®istry) {
+ registry.insert<LLVM::LLVMArmNeonDialect>();
+ registry.addDialectInterface<LLVM::LLVMArmNeonDialect,
+ LLVMArmNeonDialectLLVMIRTranslationInterface>();
+}
+
+void mlir::registerLLVMArmNeonDialectTranslation(MLIRContext &context) {
+ DialectRegistry registry;
+ registerLLVMArmNeonDialectTranslation(registry);
+ context.appendDialectRegistry(registry);
}
diff --git a/mlir/lib/Target/LLVMIR/Dialect/LLVMArmSVE/CMakeLists.txt b/mlir/lib/Target/LLVMIR/Dialect/LLVMArmSVE/CMakeLists.txt
index b62391618ba0..aab6a6c926a6 100644
--- a/mlir/lib/Target/LLVMIR/Dialect/LLVMArmSVE/CMakeLists.txt
+++ b/mlir/lib/Target/LLVMIR/Dialect/LLVMArmSVE/CMakeLists.txt
@@ -12,5 +12,5 @@ add_mlir_translation_library(MLIRLLVMArmSVEToLLVMIRTranslation
MLIRLLVMArmSVE
MLIRLLVMIR
MLIRSupport
- MLIRTargetLLVMIRModuleTranslation
+ MLIRTargetLLVMIRExport
)
diff --git a/mlir/lib/Target/LLVMIR/Dialect/LLVMArmSVE/LLVMArmSVEToLLVMIRTranslation.cpp b/mlir/lib/Target/LLVMIR/Dialect/LLVMArmSVE/LLVMArmSVEToLLVMIRTranslation.cpp
index 7bd9566b0972..72d496a925c5 100644
--- a/mlir/lib/Target/LLVMIR/Dialect/LLVMArmSVE/LLVMArmSVEToLLVMIRTranslation.cpp
+++ b/mlir/lib/Target/LLVMIR/Dialect/LLVMArmSVE/LLVMArmSVEToLLVMIRTranslation.cpp
@@ -22,12 +22,35 @@
using namespace mlir;
using namespace mlir::LLVM;
-LogicalResult
-mlir::LLVMArmSVEDialectLLVMIRTranslationInterface::convertOperation(
- Operation *op, llvm::IRBuilderBase &builder,
- LLVM::ModuleTranslation &moduleTranslation) const {
- Operation &opInst = *op;
+namespace {
+/// Implementation of the dialect interface that converts operations belonging
+/// to the LLVMArmSVE dialect to LLVM IR.
+class LLVMArmSVEDialectLLVMIRTranslationInterface
+ : public LLVMTranslationDialectInterface {
+public:
+ using LLVMTranslationDialectInterface::LLVMTranslationDialectInterface;
+
+ /// Translates the given operation to LLVM IR using the provided IR builder
+ /// and saving the state in `moduleTranslation`.
+ LogicalResult
+ convertOperation(Operation *op, llvm::IRBuilderBase &builder,
+ LLVM::ModuleTranslation &moduleTranslation) const final {
+ Operation &opInst = *op;
#include "mlir/Dialect/LLVMIR/LLVMArmSVEConversions.inc"
- return failure();
+ return failure();
+ }
+};
+} // end namespace
+
+void mlir::registerLLVMArmSVEDialectTranslation(DialectRegistry ®istry) {
+ registry.insert<LLVM::LLVMArmSVEDialect>();
+ registry.addDialectInterface<LLVM::LLVMArmSVEDialect,
+ LLVMArmSVEDialectLLVMIRTranslationInterface>();
+}
+
+void mlir::registerLLVMArmSVEDialectTranslation(MLIRContext &context) {
+ DialectRegistry registry;
+ registerLLVMArmSVEDialectTranslation(registry);
+ context.appendDialectRegistry(registry);
}
diff --git a/mlir/lib/Target/LLVMIR/Dialect/LLVMIR/CMakeLists.txt b/mlir/lib/Target/LLVMIR/Dialect/LLVMIR/CMakeLists.txt
index 2da7e95cabb5..68b89936714c 100644
--- a/mlir/lib/Target/LLVMIR/Dialect/LLVMIR/CMakeLists.txt
+++ b/mlir/lib/Target/LLVMIR/Dialect/LLVMIR/CMakeLists.txt
@@ -8,5 +8,5 @@ add_mlir_translation_library(MLIRLLVMToLLVMIRTranslation
MLIRIR
MLIRLLVMIR
MLIRSupport
- MLIRTargetLLVMIRModuleTranslation
+ MLIRTargetLLVMIRExport
)
diff --git a/mlir/lib/Target/LLVMIR/Dialect/LLVMIR/LLVMToLLVMIRTranslation.cpp b/mlir/lib/Target/LLVMIR/Dialect/LLVMIR/LLVMToLLVMIRTranslation.cpp
index 1f03092628a5..4d0e380b08ad 100644
--- a/mlir/lib/Target/LLVMIR/Dialect/LLVMIR/LLVMToLLVMIRTranslation.cpp
+++ b/mlir/lib/Target/LLVMIR/Dialect/LLVMIR/LLVMToLLVMIRTranslation.cpp
@@ -446,8 +446,32 @@ convertOperationImpl(Operation &opInst, llvm::IRBuilderBase &builder,
return failure();
}
-LogicalResult mlir::LLVMDialectLLVMIRTranslationInterface::convertOperation(
- Operation *op, llvm::IRBuilderBase &builder,
- LLVM::ModuleTranslation &moduleTranslation) const {
- return convertOperationImpl(*op, builder, moduleTranslation);
+namespace {
+/// Implementation of the dialect interface that converts operations belonging
+/// to the LLVM dialect to LLVM IR.
+class LLVMDialectLLVMIRTranslationInterface
+ : public LLVMTranslationDialectInterface {
+public:
+ using LLVMTranslationDialectInterface::LLVMTranslationDialectInterface;
+
+ /// Translates the given operation to LLVM IR using the provided IR builder
+ /// and saving the state in `moduleTranslation`.
+ LogicalResult
+ convertOperation(Operation *op, llvm::IRBuilderBase &builder,
+ LLVM::ModuleTranslation &moduleTranslation) const final {
+ return convertOperationImpl(*op, builder, moduleTranslation);
+ }
+};
+} // end namespace
+
+void mlir::registerLLVMDialectTranslation(DialectRegistry ®istry) {
+ registry.insert<LLVM::LLVMDialect>();
+ registry.addDialectInterface<LLVM::LLVMDialect,
+ LLVMDialectLLVMIRTranslationInterface>();
+}
+
+void mlir::registerLLVMDialectTranslation(MLIRContext &context) {
+ DialectRegistry registry;
+ registerLLVMDialectTranslation(registry);
+ context.appendDialectRegistry(registry);
}
diff --git a/mlir/lib/Target/LLVMIR/Dialect/NVVM/CMakeLists.txt b/mlir/lib/Target/LLVMIR/Dialect/NVVM/CMakeLists.txt
index 5a2615a84a38..4518904a110b 100644
--- a/mlir/lib/Target/LLVMIR/Dialect/NVVM/CMakeLists.txt
+++ b/mlir/lib/Target/LLVMIR/Dialect/NVVM/CMakeLists.txt
@@ -12,5 +12,5 @@ add_mlir_translation_library(MLIRNVVMToLLVMIRTranslation
MLIRLLVMIR
MLIRNVVMIR
MLIRSupport
- MLIRTargetLLVMIRModuleTranslation
+ MLIRTargetLLVMIRExport
)
diff --git a/mlir/lib/Target/LLVMIR/Dialect/NVVM/NVVMToLLVMIRTranslation.cpp b/mlir/lib/Target/LLVMIR/Dialect/NVVM/NVVMToLLVMIRTranslation.cpp
index a0eb87ec3125..bdcb451323ad 100644
--- a/mlir/lib/Target/LLVMIR/Dialect/NVVM/NVVMToLLVMIRTranslation.cpp
+++ b/mlir/lib/Target/LLVMIR/Dialect/NVVM/NVVMToLLVMIRTranslation.cpp
@@ -34,31 +34,60 @@ static llvm::Intrinsic::ID getShflBflyIntrinsicId(llvm::Type *resultType,
: llvm::Intrinsic::nvvm_shfl_sync_bfly_i32;
}
-LogicalResult mlir::NVVMDialectLLVMIRTranslationInterface::convertOperation(
- Operation *op, llvm::IRBuilderBase &builder,
- LLVM::ModuleTranslation &moduleTranslation) const {
- Operation &opInst = *op;
+namespace {
+/// Implementation of the dialect interface that converts operations belonging
+/// to the NVVM dialect to LLVM IR.
+class NVVMDialectLLVMIRTranslationInterface
+ : public LLVMTranslationDialectInterface {
+public:
+ using LLVMTranslationDialectInterface::LLVMTranslationDialectInterface;
+
+ /// Translates the given operation to LLVM IR using the provided IR builder
+ /// and saving the state in `moduleTranslation`.
+ LogicalResult
+ convertOperation(Operation *op, llvm::IRBuilderBase &builder,
+ LLVM::ModuleTranslation &moduleTranslation) const final {
+ Operation &opInst = *op;
#include "mlir/Dialect/LLVMIR/NVVMConversions.inc"
- return failure();
-}
+ return failure();
+ }
+
+ /// Attaches module-level metadata for functions marked as kernels.
+ LogicalResult
+ amendOperation(Operation *op, NamedAttribute attribute,
+ LLVM::ModuleTranslation &moduleTranslation) const final {
+ if (attribute.first == NVVM::NVVMDialect::getKernelFuncAttrName()) {
+ auto func = dyn_cast<LLVM::LLVMFuncOp>(op);
+ if (!func)
+ return failure();
-LogicalResult mlir::NVVMDialectLLVMIRTranslationInterface::amendOperation(
- Operation *op, NamedAttribute attribute,
- LLVM::ModuleTranslation &moduleTranslation) const {
- if (attribute.first == NVVM::NVVMDialect::getKernelFuncAttrName()) {
- auto func = cast<LLVM::LLVMFuncOp>(op);
- 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);
+ 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();
}
- return success();
+};
+} // end namespace
+
+void mlir::registerNVVMDialectTranslation(DialectRegistry ®istry) {
+ registry.insert<NVVM::NVVMDialect>();
+ registry.addDialectInterface<NVVM::NVVMDialect,
+ NVVMDialectLLVMIRTranslationInterface>();
+}
+
+void mlir::registerNVVMDialectTranslation(MLIRContext &context) {
+ DialectRegistry registry;
+ registerNVVMDialectTranslation(registry);
+ context.appendDialectRegistry(registry);
}
diff --git a/mlir/lib/Target/LLVMIR/Dialect/OpenMP/CMakeLists.txt b/mlir/lib/Target/LLVMIR/Dialect/OpenMP/CMakeLists.txt
index 6cc36f47f075..4a06df33e99b 100644
--- a/mlir/lib/Target/LLVMIR/Dialect/OpenMP/CMakeLists.txt
+++ b/mlir/lib/Target/LLVMIR/Dialect/OpenMP/CMakeLists.txt
@@ -9,5 +9,5 @@ add_mlir_translation_library(MLIROpenMPToLLVMIRTranslation
MLIRLLVMIR
MLIROpenMP
MLIRSupport
- MLIRTargetLLVMIRModuleTranslation
+ MLIRTargetLLVMIRExport
)
diff --git a/mlir/lib/Target/LLVMIR/Dialect/OpenMP/OpenMPToLLVMIRTranslation.cpp b/mlir/lib/Target/LLVMIR/Dialect/OpenMP/OpenMPToLLVMIRTranslation.cpp
index 9a2bee9482e5..fbb50ea2eea8 100644
--- a/mlir/lib/Target/LLVMIR/Dialect/OpenMP/OpenMPToLLVMIRTranslation.cpp
+++ b/mlir/lib/Target/LLVMIR/Dialect/OpenMP/OpenMPToLLVMIRTranslation.cpp
@@ -168,8 +168,9 @@ convertOmpMaster(Operation &opInst, llvm::IRBuilderBase &builder,
}
/// Converts an OpenMP workshare loop into LLVM IR using OpenMPIRBuilder.
-LogicalResult convertOmpWsLoop(Operation &opInst, llvm::IRBuilderBase &builder,
- LLVM::ModuleTranslation &moduleTranslation) {
+static LogicalResult
+convertOmpWsLoop(Operation &opInst, llvm::IRBuilderBase &builder,
+ LLVM::ModuleTranslation &moduleTranslation) {
auto loop = cast<omp::WsLoopOp>(opInst);
// TODO: this should be in the op verifier instead.
if (loop.lowerBound().empty())
@@ -248,9 +249,27 @@ LogicalResult convertOmpWsLoop(Operation &opInst, llvm::IRBuilderBase &builder,
return success();
}
+namespace {
+
+/// Implementation of the dialect interface that converts operations belonging
+/// to the OpenMP dialect to LLVM IR.
+class OpenMPDialectLLVMIRTranslationInterface
+ : public LLVMTranslationDialectInterface {
+public:
+ using LLVMTranslationDialectInterface::LLVMTranslationDialectInterface;
+
+ /// Translates the given operation to LLVM IR using the provided IR builder
+ /// and saving the state in `moduleTranslation`.
+ LogicalResult
+ convertOperation(Operation *op, llvm::IRBuilderBase &builder,
+ LLVM::ModuleTranslation &moduleTranslation) const final;
+};
+
+} // end namespace
+
/// Given an OpenMP MLIR operation, create the corresponding LLVM IR
/// (including OpenMP runtime calls).
-LogicalResult mlir::OpenMPDialectLLVMIRTranslationInterface::convertOperation(
+LogicalResult OpenMPDialectLLVMIRTranslationInterface::convertOperation(
Operation *op, llvm::IRBuilderBase &builder,
LLVM::ModuleTranslation &moduleTranslation) const {
@@ -302,3 +321,15 @@ LogicalResult mlir::OpenMPDialectLLVMIRTranslationInterface::convertOperation(
<< inst->getName();
});
}
+
+void mlir::registerOpenMPDialectTranslation(DialectRegistry ®istry) {
+ registry.insert<omp::OpenMPDialect>();
+ registry.addDialectInterface<omp::OpenMPDialect,
+ OpenMPDialectLLVMIRTranslationInterface>();
+}
+
+void mlir::registerOpenMPDialectTranslation(MLIRContext &context) {
+ DialectRegistry registry;
+ registerOpenMPDialectTranslation(registry);
+ context.appendDialectRegistry(registry);
+}
diff --git a/mlir/lib/Target/LLVMIR/Dialect/ROCDL/CMakeLists.txt b/mlir/lib/Target/LLVMIR/Dialect/ROCDL/CMakeLists.txt
index 7e542ec77adb..5b6323d5e486 100644
--- a/mlir/lib/Target/LLVMIR/Dialect/ROCDL/CMakeLists.txt
+++ b/mlir/lib/Target/LLVMIR/Dialect/ROCDL/CMakeLists.txt
@@ -12,5 +12,5 @@ add_mlir_translation_library(MLIRROCDLToLLVMIRTranslation
MLIRLLVMIR
MLIRROCDLIR
MLIRSupport
- MLIRTargetLLVMIRModuleTranslation
+ MLIRTargetLLVMIRExport
)
diff --git a/mlir/lib/Target/LLVMIR/Dialect/ROCDL/ROCDLToLLVMIRTranslation.cpp b/mlir/lib/Target/LLVMIR/Dialect/ROCDL/ROCDLToLLVMIRTranslation.cpp
index 9288af17beef..8a3470ce80aa 100644
--- a/mlir/lib/Target/LLVMIR/Dialect/ROCDL/ROCDLToLLVMIRTranslation.cpp
+++ b/mlir/lib/Target/LLVMIR/Dialect/ROCDL/ROCDLToLLVMIRTranslation.cpp
@@ -41,27 +41,55 @@ static llvm::Value *createDeviceFunctionCall(llvm::IRBuilderBase &builder,
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;
+namespace {
+/// Implementation of the dialect interface that converts operations belonging
+/// to the ROCDL dialect to LLVM IR.
+class ROCDLDialectLLVMIRTranslationInterface
+ : public LLVMTranslationDialectInterface {
+public:
+ using LLVMTranslationDialectInterface::LLVMTranslationDialectInterface;
+
+ /// Translates the given operation to LLVM IR using the provided IR builder
+ /// and saving the state in `moduleTranslation`.
+ LogicalResult
+ convertOperation(Operation *op, llvm::IRBuilderBase &builder,
+ LLVM::ModuleTranslation &moduleTranslation) const final {
+ Operation &opInst = *op;
#include "mlir/Dialect/LLVMIR/ROCDLConversions.inc"
- return failure();
-}
+ return failure();
+ }
-LogicalResult mlir::ROCDLDialectLLVMIRTranslationInterface::amendOperation(
- Operation *op, NamedAttribute attribute,
- LLVM::ModuleTranslation &moduleTranslation) const {
- if (attribute.first == ROCDL::ROCDLDialect::getKernelFuncAttrName()) {
- auto func = cast<LLVM::LLVMFuncOp>(op);
+ /// Attaches module-level metadata for functions marked as kernels.
+ LogicalResult
+ amendOperation(Operation *op, NamedAttribute attribute,
+ LLVM::ModuleTranslation &moduleTranslation) const final {
+ if (attribute.first == ROCDL::ROCDLDialect::getKernelFuncAttrName()) {
+ auto func = 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");
+ // 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();
}
- return success();
+};
+} // end namespace
+
+void mlir::registerROCDLDialectTranslation(DialectRegistry ®istry) {
+ registry.insert<ROCDL::ROCDLDialect>();
+ registry.addDialectInterface<ROCDL::ROCDLDialect,
+ ROCDLDialectLLVMIRTranslationInterface>();
+}
+
+void mlir::registerROCDLDialectTranslation(MLIRContext &context) {
+ DialectRegistry registry;
+ registerROCDLDialectTranslation(registry);
+ context.appendDialectRegistry(registry);
}
diff --git a/mlir/lib/Target/LLVMIR/TypeTranslation.cpp b/mlir/lib/Target/LLVMIR/TypeTranslation.cpp
index 7a1ff099bb86..f40a6d943cbb 100644
--- a/mlir/lib/Target/LLVMIR/TypeTranslation.cpp
+++ b/mlir/lib/Target/LLVMIR/TypeTranslation.cpp
@@ -187,147 +187,3 @@ unsigned LLVM::TypeToLLVMIRTranslator::getPreferredAlignment(
Type type, const llvm::DataLayout &layout) {
return layout.getPrefTypeAlignment(translateType(type));
}
-
-namespace mlir {
-namespace LLVM {
-namespace detail {
-/// Support for translating LLVM IR types to MLIR LLVM dialect types.
-class TypeFromLLVMIRTranslatorImpl {
-public:
- /// Constructs a class creating types in the given MLIR context.
- TypeFromLLVMIRTranslatorImpl(MLIRContext &context) : context(context) {}
-
- /// Translates the given type.
- Type translateType(llvm::Type *type) {
- if (knownTranslations.count(type))
- return knownTranslations.lookup(type);
-
- Type translated =
- llvm::TypeSwitch<llvm::Type *, Type>(type)
- .Case<llvm::ArrayType, llvm::FunctionType, llvm::IntegerType,
- llvm::PointerType, llvm::StructType, llvm::FixedVectorType,
- llvm::ScalableVectorType>(
- [this](auto *type) { return this->translate(type); })
- .Default([this](llvm::Type *type) {
- return translatePrimitiveType(type);
- });
- knownTranslations.try_emplace(type, translated);
- return translated;
- }
-
-private:
- /// Translates the given primitive, i.e. non-parametric in MLIR nomenclature,
- /// type.
- Type translatePrimitiveType(llvm::Type *type) {
- if (type->isVoidTy())
- return LLVM::LLVMVoidType::get(&context);
- if (type->isHalfTy())
- return Float16Type::get(&context);
- if (type->isBFloatTy())
- return BFloat16Type::get(&context);
- if (type->isFloatTy())
- return Float32Type::get(&context);
- if (type->isDoubleTy())
- return Float64Type::get(&context);
- if (type->isFP128Ty())
- return Float128Type::get(&context);
- if (type->isX86_FP80Ty())
- return Float80Type::get(&context);
- if (type->isPPC_FP128Ty())
- return LLVM::LLVMPPCFP128Type::get(&context);
- if (type->isX86_MMXTy())
- return LLVM::LLVMX86MMXType::get(&context);
- if (type->isLabelTy())
- return LLVM::LLVMLabelType::get(&context);
- if (type->isMetadataTy())
- return LLVM::LLVMMetadataType::get(&context);
- llvm_unreachable("not a primitive type");
- }
-
- /// Translates the given array type.
- Type translate(llvm::ArrayType *type) {
- return LLVM::LLVMArrayType::get(translateType(type->getElementType()),
- type->getNumElements());
- }
-
- /// Translates the given function type.
- Type translate(llvm::FunctionType *type) {
- SmallVector<Type, 8> paramTypes;
- translateTypes(type->params(), paramTypes);
- return LLVM::LLVMFunctionType::get(translateType(type->getReturnType()),
- paramTypes, type->isVarArg());
- }
-
- /// Translates the given integer type.
- Type translate(llvm::IntegerType *type) {
- return IntegerType::get(&context, type->getBitWidth());
- }
-
- /// Translates the given pointer type.
- Type translate(llvm::PointerType *type) {
- return LLVM::LLVMPointerType::get(translateType(type->getElementType()),
- type->getAddressSpace());
- }
-
- /// Translates the given structure type.
- Type translate(llvm::StructType *type) {
- SmallVector<Type, 8> subtypes;
- if (type->isLiteral()) {
- translateTypes(type->subtypes(), subtypes);
- return LLVM::LLVMStructType::getLiteral(&context, subtypes,
- type->isPacked());
- }
-
- if (type->isOpaque())
- return LLVM::LLVMStructType::getOpaque(type->getName(), &context);
-
- LLVM::LLVMStructType translated =
- LLVM::LLVMStructType::getIdentified(&context, type->getName());
- knownTranslations.try_emplace(type, translated);
- translateTypes(type->subtypes(), subtypes);
- LogicalResult bodySet = translated.setBody(subtypes, type->isPacked());
- assert(succeeded(bodySet) &&
- "could not set the body of an identified struct");
- (void)bodySet;
- return translated;
- }
-
- /// Translates the given fixed-vector type.
- Type translate(llvm::FixedVectorType *type) {
- return LLVM::getFixedVectorType(translateType(type->getElementType()),
- type->getNumElements());
- }
-
- /// Translates the given scalable-vector type.
- Type translate(llvm::ScalableVectorType *type) {
- return LLVM::LLVMScalableVectorType::get(
- translateType(type->getElementType()), type->getMinNumElements());
- }
-
- /// Translates a list of types.
- void translateTypes(ArrayRef<llvm::Type *> types,
- SmallVectorImpl<Type> &result) {
- result.reserve(result.size() + types.size());
- for (llvm::Type *type : types)
- result.push_back(translateType(type));
- }
-
- /// Map of known translations. Serves as a cache and as recursion stopper for
- /// translating recursive structs.
- llvm::DenseMap<llvm::Type *, Type> knownTranslations;
-
- /// The context in which MLIR types are created.
- MLIRContext &context;
-};
-} // end namespace detail
-} // end namespace LLVM
-} // end namespace mlir
-
-LLVM::TypeFromLLVMIRTranslator::TypeFromLLVMIRTranslator(MLIRContext &context)
- : impl(new detail::TypeFromLLVMIRTranslatorImpl(context)) {}
-
-LLVM::TypeFromLLVMIRTranslator::~TypeFromLLVMIRTranslator() {}
-
-Type LLVM::TypeFromLLVMIRTranslator::translateType(llvm::Type *type) {
- return impl->translateType(type);
-}
diff --git a/mlir/test/lib/Transforms/CMakeLists.txt b/mlir/test/lib/Transforms/CMakeLists.txt
index 7de188df4aa8..b5ae9ecef0de 100644
--- a/mlir/test/lib/Transforms/CMakeLists.txt
+++ b/mlir/test/lib/Transforms/CMakeLists.txt
@@ -51,6 +51,7 @@ add_mlir_library(MLIRTestTransforms
MLIRGPUToGPURuntimeTransforms
MLIRLinalg
MLIRLinalgTransforms
+ MLIRLLVMToLLVMIRTranslation
MLIRMathTransforms
MLIRNVVMIR
MLIRNVVMToLLVMIRTranslation
@@ -60,7 +61,7 @@ add_mlir_library(MLIRTestTransforms
MLIRSCF
MLIRSCFTransforms
MLIRStandardOpsTransforms
- MLIRTargetLLVMIR
+ MLIRTargetLLVMIRExport
MLIRTestDialect
MLIRTransformUtils
MLIRVector
diff --git a/mlir/test/lib/Transforms/TestConvertGPUKernelToCubin.cpp b/mlir/test/lib/Transforms/TestConvertGPUKernelToCubin.cpp
index ab6448b8a95e..0f8c4959f2cc 100644
--- a/mlir/test/lib/Transforms/TestConvertGPUKernelToCubin.cpp
+++ b/mlir/test/lib/Transforms/TestConvertGPUKernelToCubin.cpp
@@ -10,7 +10,7 @@
#include "mlir/Dialect/LLVMIR/NVVMDialect.h"
#include "mlir/Pass/Pass.h"
#include "mlir/Pass/PassManager.h"
-#include "mlir/Target/LLVMIR.h"
+#include "mlir/Target/LLVMIR/Dialect/LLVMIR/LLVMToLLVMIRTranslation.h"
#include "mlir/Target/LLVMIR/Dialect/NVVM/NVVMToLLVMIRTranslation.h"
#include "mlir/Target/LLVMIR/Export.h"
#include "llvm/Support/TargetSelect.h"
@@ -24,19 +24,6 @@ static OwnedBlob compilePtxToCubinForTesting(const std::string &, Location,
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) {
diff --git a/mlir/test/lib/Transforms/TestConvertGPUKernelToHsaco.cpp b/mlir/test/lib/Transforms/TestConvertGPUKernelToHsaco.cpp
index 49db2f28ede3..58e890b907a4 100644
--- a/mlir/test/lib/Transforms/TestConvertGPUKernelToHsaco.cpp
+++ b/mlir/test/lib/Transforms/TestConvertGPUKernelToHsaco.cpp
@@ -10,7 +10,7 @@
#include "mlir/Dialect/LLVMIR/ROCDLDialect.h"
#include "mlir/Pass/Pass.h"
#include "mlir/Pass/PassManager.h"
-#include "mlir/Target/LLVMIR.h"
+#include "mlir/Target/LLVMIR/Dialect/LLVMIR/LLVMToLLVMIRTranslation.h"
#include "mlir/Target/LLVMIR/Dialect/ROCDL/ROCDLToLLVMIRTranslation.h"
#include "mlir/Target/LLVMIR/Export.h"
#include "llvm/Support/TargetSelect.h"
@@ -24,14 +24,6 @@ static OwnedBlob compileIsaToHsacoForTesting(const std::string &, Location,
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) {
diff --git a/mlir/tools/mlir-cpu-runner/CMakeLists.txt b/mlir/tools/mlir-cpu-runner/CMakeLists.txt
index 8ff73850ee68..1135a683eb6e 100644
--- a/mlir/tools/mlir-cpu-runner/CMakeLists.txt
+++ b/mlir/tools/mlir-cpu-runner/CMakeLists.txt
@@ -16,8 +16,10 @@ target_link_libraries(mlir-cpu-runner PRIVATE
MLIRIR
MLIRJitRunner
MLIRLLVMIR
+ MLIRLLVMToLLVMIRTranslation
MLIROpenMP
+ MLIROpenMPToLLVMIRTranslation
MLIRParser
- MLIRTargetLLVMIR
+ MLIRTargetLLVMIRExport
MLIRSupport
)
diff --git a/mlir/tools/mlir-cpu-runner/mlir-cpu-runner.cpp b/mlir/tools/mlir-cpu-runner/mlir-cpu-runner.cpp
index 4ce48782672f..d22b3a20bc8b 100644
--- a/mlir/tools/mlir-cpu-runner/mlir-cpu-runner.cpp
+++ b/mlir/tools/mlir-cpu-runner/mlir-cpu-runner.cpp
@@ -13,11 +13,10 @@
//===----------------------------------------------------------------------===//
#include "mlir/Dialect/LLVMIR/LLVMDialect.h"
-#include "mlir/Dialect/OpenMP/OpenMPDialect.h"
#include "mlir/ExecutionEngine/JitRunner.h"
#include "mlir/ExecutionEngine/OptUtils.h"
#include "mlir/IR/Dialect.h"
-#include "mlir/Target/LLVMIR.h"
+#include "mlir/Target/LLVMIR/Dialect/LLVMIR/LLVMToLLVMIRTranslation.h"
#include "mlir/Target/LLVMIR/Dialect/OpenMP/OpenMPToLLVMIRTranslation.h"
#include "llvm/Support/InitLLVM.h"
@@ -31,10 +30,8 @@ int main(int argc, char **argv) {
mlir::initializeLLVMPasses();
mlir::DialectRegistry registry;
- registry.insert<mlir::LLVM::LLVMDialect, mlir::omp::OpenMPDialect>();
mlir::registerLLVMDialectTranslation(registry);
- registry.addDialectInterface<mlir::omp::OpenMPDialect,
- mlir::OpenMPDialectLLVMIRTranslationInterface>();
+ mlir::registerOpenMPDialectTranslation(registry);
return mlir::JitRunnerMain(argc, argv, registry);
}
diff --git a/mlir/tools/mlir-cuda-runner/CMakeLists.txt b/mlir/tools/mlir-cuda-runner/CMakeLists.txt
index 62ded09bc789..be585b43b29e 100644
--- a/mlir/tools/mlir-cuda-runner/CMakeLists.txt
+++ b/mlir/tools/mlir-cuda-runner/CMakeLists.txt
@@ -37,11 +37,12 @@ if(MLIR_CUDA_RUNNER_ENABLED)
MLIRGPU
MLIRIR
MLIRLLVMIR
+ MLIRLLVMToLLVMIRTranslation
MLIRNVVMIR
MLIRParser
MLIRStandard
MLIRSupport
- MLIRTargetLLVMIR
+ MLIRTargetLLVMIRExport
MLIRNVVMToLLVMIRTranslation
MLIRTransforms
MLIRTranslation
diff --git a/mlir/tools/mlir-cuda-runner/mlir-cuda-runner.cpp b/mlir/tools/mlir-cuda-runner/mlir-cuda-runner.cpp
index 025b8133d8d9..fde875927f1f 100644
--- a/mlir/tools/mlir-cuda-runner/mlir-cuda-runner.cpp
+++ b/mlir/tools/mlir-cuda-runner/mlir-cuda-runner.cpp
@@ -29,7 +29,7 @@
#include "mlir/IR/BuiltinOps.h"
#include "mlir/Pass/Pass.h"
#include "mlir/Pass/PassManager.h"
-#include "mlir/Target/LLVMIR.h"
+#include "mlir/Target/LLVMIR/Dialect/LLVMIR/LLVMToLLVMIRTranslation.h"
#include "mlir/Target/LLVMIR/Dialect/NVVM/NVVMToLLVMIRTranslation.h"
#include "mlir/Target/LLVMIR/Export.h"
#include "mlir/Transforms/DialectConversion.h"
@@ -181,9 +181,8 @@ int main(int argc, char **argv) {
registry.insert<mlir::LLVM::LLVMDialect, mlir::NVVM::NVVMDialect,
mlir::async::AsyncDialect, mlir::gpu::GPUDialect,
mlir::StandardOpsDialect>();
- registry.addDialectInterface<NVVM::NVVMDialect,
- NVVMDialectLLVMIRTranslationInterface>();
mlir::registerLLVMDialectTranslation(registry);
+ mlir::registerNVVMDialectTranslation(registry);
return mlir::JitRunnerMain(argc, argv, registry, jitRunnerConfig);
}
diff --git a/mlir/tools/mlir-rocm-runner/CMakeLists.txt b/mlir/tools/mlir-rocm-runner/CMakeLists.txt
index 6b649d29edb5..d2381413e158 100644
--- a/mlir/tools/mlir-rocm-runner/CMakeLists.txt
+++ b/mlir/tools/mlir-rocm-runner/CMakeLists.txt
@@ -78,11 +78,12 @@ if(MLIR_ROCM_RUNNER_ENABLED)
MLIRGPU
MLIRIR
MLIRLLVMIR
+ MLIRLLVMToLLVMIRTranslation
MLIRParser
MLIRROCDLIR
MLIRStandard
MLIRSupport
- MLIRTargetLLVMIR
+ MLIRTargetLLVMIRExport
MLIRROCDLToLLVMIRTranslation
MLIRTransforms
MLIRTranslation
diff --git a/mlir/tools/mlir-rocm-runner/mlir-rocm-runner.cpp b/mlir/tools/mlir-rocm-runner/mlir-rocm-runner.cpp
index ac579da9728c..c2f9abbf73c8 100644
--- a/mlir/tools/mlir-rocm-runner/mlir-rocm-runner.cpp
+++ b/mlir/tools/mlir-rocm-runner/mlir-rocm-runner.cpp
@@ -30,7 +30,7 @@
#include "mlir/Pass/Pass.h"
#include "mlir/Pass/PassManager.h"
#include "mlir/Support/FileUtilities.h"
-#include "mlir/Target/LLVMIR.h"
+#include "mlir/Target/LLVMIR/Dialect/LLVMIR/LLVMToLLVMIRTranslation.h"
#include "mlir/Target/LLVMIR/Dialect/ROCDL/ROCDLToLLVMIRTranslation.h"
#include "mlir/Target/LLVMIR/Export.h"
#include "mlir/Transforms/DialectConversion.h"
@@ -342,9 +342,8 @@ int main(int argc, char **argv) {
mlir::DialectRegistry registry;
registry.insert<mlir::LLVM::LLVMDialect, mlir::gpu::GPUDialect,
mlir::ROCDL::ROCDLDialect, mlir::StandardOpsDialect>();
- registry.addDialectInterface<ROCDL::ROCDLDialect,
- ROCDLDialectLLVMIRTranslationInterface>();
mlir::registerLLVMDialectTranslation(registry);
+ mlir::registerROCDLDialectTranslation(registry);
return mlir::JitRunnerMain(argc, argv, registry, jitRunnerConfig);
}
diff --git a/mlir/tools/mlir-spirv-cpu-runner/CMakeLists.txt b/mlir/tools/mlir-spirv-cpu-runner/CMakeLists.txt
index e7b8217958c3..685bddcfc6a5 100644
--- a/mlir/tools/mlir-spirv-cpu-runner/CMakeLists.txt
+++ b/mlir/tools/mlir-spirv-cpu-runner/CMakeLists.txt
@@ -22,10 +22,11 @@ if (MLIR_SPIRV_CPU_RUNNER_ENABLED)
MLIRIR
MLIRJitRunner
MLIRLLVMIR
+ MLIRLLVMToLLVMIRTranslation
MLIRParser
MLIRSPIRV
MLIRStandard
- MLIRTargetLLVMIR
+ MLIRTargetLLVMIRExport
MLIRTransforms
MLIRTranslation
MLIRSupport
diff --git a/mlir/tools/mlir-spirv-cpu-runner/mlir-spirv-cpu-runner.cpp b/mlir/tools/mlir-spirv-cpu-runner/mlir-spirv-cpu-runner.cpp
index cc098e79dc77..0b1b2c4a52eb 100644
--- a/mlir/tools/mlir-spirv-cpu-runner/mlir-spirv-cpu-runner.cpp
+++ b/mlir/tools/mlir-spirv-cpu-runner/mlir-spirv-cpu-runner.cpp
@@ -26,7 +26,7 @@
#include "mlir/ExecutionEngine/OptUtils.h"
#include "mlir/Pass/Pass.h"
#include "mlir/Pass/PassManager.h"
-#include "mlir/Target/LLVMIR.h"
+#include "mlir/Target/LLVMIR/Dialect/LLVMIR/LLVMToLLVMIRTranslation.h"
#include "mlir/Target/LLVMIR/Export.h"
#include "llvm/IR/LLVMContext.h"
diff --git a/mlir/tools/mlir-vulkan-runner/CMakeLists.txt b/mlir/tools/mlir-vulkan-runner/CMakeLists.txt
index 397ccc8ac707..3414456b9085 100644
--- a/mlir/tools/mlir-vulkan-runner/CMakeLists.txt
+++ b/mlir/tools/mlir-vulkan-runner/CMakeLists.txt
@@ -62,12 +62,13 @@ if (MLIR_VULKAN_RUNNER_ENABLED)
MLIRIR
MLIRJitRunner
MLIRLLVMIR
+ MLIRLLVMToLLVMIRTranslation
MLIRParser
MLIRSPIRV
MLIRSPIRVTransforms
MLIRStandard
MLIRSupport
- MLIRTargetLLVMIR
+ MLIRTargetLLVMIRExport
MLIRTransforms
MLIRTranslation
${Vulkan_LIBRARY}
diff --git a/mlir/tools/mlir-vulkan-runner/mlir-vulkan-runner.cpp b/mlir/tools/mlir-vulkan-runner/mlir-vulkan-runner.cpp
index 651e59cddf4f..af74b5d3057b 100644
--- a/mlir/tools/mlir-vulkan-runner/mlir-vulkan-runner.cpp
+++ b/mlir/tools/mlir-vulkan-runner/mlir-vulkan-runner.cpp
@@ -27,7 +27,8 @@
#include "mlir/ExecutionEngine/OptUtils.h"
#include "mlir/Pass/Pass.h"
#include "mlir/Pass/PassManager.h"
-#include "mlir/Target/LLVMIR.h"
+#include "mlir/Target/LLVMIR/Dialect/LLVMIR/LLVMToLLVMIRTranslation.h"
+#include "mlir/Target/LLVMIR/Export.h"
#include "llvm/Support/InitLLVM.h"
#include "llvm/Support/TargetSelect.h"
diff --git a/mlir/unittests/ExecutionEngine/Invoke.cpp b/mlir/unittests/ExecutionEngine/Invoke.cpp
index 230c9b6626e6..f970cc221985 100644
--- a/mlir/unittests/ExecutionEngine/Invoke.cpp
+++ b/mlir/unittests/ExecutionEngine/Invoke.cpp
@@ -19,7 +19,8 @@
#include "mlir/InitAllDialects.h"
#include "mlir/Parser.h"
#include "mlir/Pass/PassManager.h"
-#include "mlir/Target/LLVMIR.h"
+#include "mlir/Target/LLVMIR/Dialect/LLVMIR/LLVMToLLVMIRTranslation.h"
+#include "mlir/Target/LLVMIR/Export.h"
#include "llvm/Support/TargetSelect.h"
#include "llvm/Support/raw_ostream.h"
@@ -231,7 +232,7 @@ TEST(NativeMemRefJit, JITCallback) {
elt = count++;
std::string moduleStr = R"mlir(
- func private @callback(%arg0: memref<?x?xf32>, %coefficient: i32) attributes { llvm.emit_c_interface }
+ func private @callback(%arg0: memref<?x?xf32>, %coefficient: i32) attributes { llvm.emit_c_interface }
func @caller_for_callback(%arg0: memref<?x?xf32>, %coefficient: i32) attributes { llvm.emit_c_interface } {
%unranked = memref_cast %arg0: memref<?x?xf32> to memref<*xf32>
call @callback(%arg0, %coefficient) : (memref<?x?xf32>, i32) -> ()
More information about the Mlir-commits
mailing list