[Mlir-commits] [mlir] 176379e - [mlir] Use the interface-based translation for LLVM "intrinsic" dialects
Alex Zinenko
llvmlistbot at llvm.org
Mon Feb 15 05:43:18 PST 2021
Author: Alex Zinenko
Date: 2021-02-15T14:43:07+01:00
New Revision: 176379e0c8f9dbde2b357fb3b6a6802b83282e71
URL: https://github.com/llvm/llvm-project/commit/176379e0c8f9dbde2b357fb3b6a6802b83282e71
DIFF: https://github.com/llvm/llvm-project/commit/176379e0c8f9dbde2b357fb3b6a6802b83282e71.diff
LOG: [mlir] Use the interface-based translation for LLVM "intrinsic" dialects
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
Reviewed By: nicolasvasilache
Differential Revision: https://reviews.llvm.org/D96592
Added:
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/NVVM/NVVMToLLVMIRTranslation.h
mlir/include/mlir/Target/LLVMIR/Dialect/ROCDL/ROCDLToLLVMIRTranslation.h
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/NVVM/CMakeLists.txt
mlir/lib/Target/LLVMIR/Dialect/NVVM/NVVMToLLVMIRTranslation.cpp
mlir/lib/Target/LLVMIR/Dialect/ROCDL/CMakeLists.txt
mlir/lib/Target/LLVMIR/Dialect/ROCDL/ROCDLToLLVMIRTranslation.cpp
Modified:
mlir/examples/standalone/test/Standalone/standalone-translate.mlir
mlir/include/mlir/Dialect/LLVMIR/LLVMOpBase.td
mlir/include/mlir/Dialect/LLVMIR/LLVMOps.td
mlir/include/mlir/Dialect/LLVMIR/ROCDLOps.td
mlir/include/mlir/InitAllTranslations.h
mlir/include/mlir/Target/LLVMIR.h
mlir/include/mlir/Target/LLVMIR/Dialect/LLVMIR/LLVMToLLVMIRTranslation.h
mlir/include/mlir/Target/LLVMIR/Dialect/OpenMP/OpenMPToLLVMIRTranslation.h
mlir/include/mlir/Target/LLVMIR/LLVMTranslationInterface.h
mlir/include/mlir/Target/LLVMIR/ModuleTranslation.h
mlir/lib/Target/CMakeLists.txt
mlir/lib/Target/LLVMIR/ConvertToLLVMIR.cpp
mlir/lib/Target/LLVMIR/Dialect/CMakeLists.txt
mlir/lib/Target/LLVMIR/Dialect/LLVMIR/LLVMToLLVMIRTranslation.cpp
mlir/lib/Target/LLVMIR/ModuleTranslation.cpp
mlir/test/Target/arm-neon.mlir
mlir/test/Target/arm-sve.mlir
mlir/test/Target/avx512.mlir
mlir/test/Target/nvvmir.mlir
mlir/test/Target/rocdl.mlir
mlir/test/lib/Transforms/CMakeLists.txt
mlir/test/lib/Transforms/TestConvertGPUKernelToCubin.cpp
mlir/test/lib/Transforms/TestConvertGPUKernelToHsaco.cpp
mlir/tools/mlir-cuda-runner/mlir-cuda-runner.cpp
mlir/tools/mlir-rocm-runner/mlir-rocm-runner.cpp
mlir/tools/mlir-tblgen/LLVMIRConversionGen.cpp
Removed:
mlir/include/mlir/Target/NVVMIR.h
mlir/include/mlir/Target/ROCDLIR.h
mlir/lib/Target/LLVMIR/ConvertToNVVMIR.cpp
mlir/lib/Target/LLVMIR/ConvertToROCDLIR.cpp
mlir/lib/Target/LLVMIR/LLVMAVX512Intr.cpp
mlir/lib/Target/LLVMIR/LLVMArmNeonIntr.cpp
mlir/lib/Target/LLVMIR/LLVMArmSVEIntr.cpp
################################################################################
diff --git a/mlir/examples/standalone/test/Standalone/standalone-translate.mlir b/mlir/examples/standalone/test/Standalone/standalone-translate.mlir
index 2a096c38e128..16d49785ee16 100644
--- a/mlir/examples/standalone/test/Standalone/standalone-translate.mlir
+++ b/mlir/examples/standalone/test/Standalone/standalone-translate.mlir
@@ -1,8 +1,5 @@
// 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
diff --git a/mlir/include/mlir/Dialect/LLVMIR/LLVMOpBase.td b/mlir/include/mlir/Dialect/LLVMIR/LLVMOpBase.td
index ad886e55b4e6..541f7ebfadfa 100644
--- a/mlir/include/mlir/Dialect/LLVMIR/LLVMOpBase.td
+++ b/mlir/include/mlir/Dialect/LLVMIR/LLVMOpBase.td
@@ -219,12 +219,13 @@ class ListIntSubst<string pattern, list<int> values> {
// 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>()
+ .getBody()[$0])}];
}
@@ -259,7 +260,7 @@ class LLVM_IntrOpBase<Dialect dialect, string opName, string enumName,
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);
}];
@@ -325,7 +326,7 @@ class LLVM_VectorReductionAcc<string 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);
diff --git a/mlir/include/mlir/Dialect/LLVMIR/LLVMOps.td b/mlir/include/mlir/Dialect/LLVMIR/LLVMOps.td
index 7a2152b9a481..7dca08a43f7a 100644
--- a/mlir/include/mlir/Dialect/LLVMIR/LLVMOps.td
+++ b/mlir/include/mlir/Dialect/LLVMIR/LLVMOps.td
@@ -1083,7 +1083,8 @@ def LLVM_UndefOp : LLVM_Op<"mlir.undef", [NoSideEffect]>,
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 = [{
diff --git a/mlir/include/mlir/Dialect/LLVMIR/ROCDLOps.td b/mlir/include/mlir/Dialect/LLVMIR/ROCDLOps.td
index d5eec3cebb4a..cfb08ff465a2 100644
--- a/mlir/include/mlir/Dialect/LLVMIR/ROCDLOps.td
+++ b/mlir/include/mlir/Dialect/LLVMIR/ROCDLOps.td
@@ -175,7 +175,7 @@ def ROCDL_MubufStoreOp :
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});
diff --git a/mlir/include/mlir/InitAllTranslations.h b/mlir/include/mlir/InitAllTranslations.h
index 16dd113d14cd..fc319c09a8c8 100644
--- a/mlir/include/mlir/InitAllTranslations.h
+++ b/mlir/include/mlir/InitAllTranslations.h
@@ -20,11 +20,6 @@ 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
@@ -35,11 +30,6 @@ inline void registerAllTranslations() {
registerFromSPIRVTranslation();
registerToLLVMIRTranslation();
registerToSPIRVTranslation();
- registerToNVVMIRTranslation();
- registerToROCDLIRTranslation();
- registerArmNeonToLLVMIRTranslation();
- registerAVX512ToLLVMIRTranslation();
- registerArmSVEToLLVMIRTranslation();
return true;
}();
(void)initOnce;
diff --git a/mlir/include/mlir/Target/LLVMIR.h b/mlir/include/mlir/Target/LLVMIR.h
index 2050c63df73f..10bec79f3506 100644
--- a/mlir/include/mlir/Target/LLVMIR.h
+++ b/mlir/include/mlir/Target/LLVMIR.h
@@ -28,14 +28,14 @@ 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
diff --git a/mlir/include/mlir/Target/LLVMIR/Dialect/LLVMAVX512/LLVMAVX512ToLLVMIRTranslation.h b/mlir/include/mlir/Target/LLVMIR/Dialect/LLVMAVX512/LLVMAVX512ToLLVMIRTranslation.h
new file mode 100644
index 000000000000..e591a95f0357
--- /dev/null
+++ b/mlir/include/mlir/Target/LLVMIR/Dialect/LLVMAVX512/LLVMAVX512ToLLVMIRTranslation.h
@@ -0,0 +1,37 @@
+//===- LLVMAVX512ToLLVMIRTranslation.h - LLVMAVX512 to LLVM IR --*- C++ -*-===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+//
+// This file implements the dialect interface for translating the LLVMAVX512
+// dialect to LLVM IR.
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef MLIR_TARGET_LLVMIR_DIALECT_LLVMAVX512_LLVMAVX512TOLLVMIRTRANSLATION_H
+#define MLIR_TARGET_LLVMIR_DIALECT_LLVMAVX512_LLVMAVX512TOLLVMIRTRANSLATION_H
+
+#include "mlir/Target/LLVMIR/LLVMTranslationInterface.h"
+
+namespace mlir {
+
+/// Implementation of the dialect interface that converts operations belonging
+/// to the LLVMAVX512 dialect to LLVM IR.
+class LLVMAVX512DialectLLVMIRTranslationInterface
+ : public LLVMTranslationDialectInterface {
+public:
+ using LLVMTranslationDialectInterface::LLVMTranslationDialectInterface;
+
+ /// Translates the given operation to LLVM IR using the provided IR builder
+ /// and saving the state in `moduleTranslation`.
+ LogicalResult
+ convertOperation(Operation *op, llvm::IRBuilderBase &builder,
+ LLVM::ModuleTranslation &moduleTranslation) const final;
+};
+
+} // namespace mlir
+
+#endif // MLIR_TARGET_LLVMIR_DIALECT_LLVMAVX512_LLVMAVX512TOLLVMIRTRANSLATION_H
diff --git a/mlir/include/mlir/Target/LLVMIR/Dialect/LLVMArmNeon/LLVMArmNeonToLLVMIRTranslation.h b/mlir/include/mlir/Target/LLVMIR/Dialect/LLVMArmNeon/LLVMArmNeonToLLVMIRTranslation.h
new file mode 100644
index 000000000000..7d268d155083
--- /dev/null
+++ b/mlir/include/mlir/Target/LLVMIR/Dialect/LLVMArmNeon/LLVMArmNeonToLLVMIRTranslation.h
@@ -0,0 +1,37 @@
+//===- LLVMArmNeonToLLVMIRTranslation.h - LLVMArmNeon to LLVMIR -*- C++ -*-===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+//
+// This file implements the dialect interface for translating the LLVMArmNeon
+// dialect to LLVM IR.
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef MLIR_TARGET_LLVMIR_DIALECT_LLVMARMNEON_LLVMARMNEONTOLLVMIRTRANSLATION_H
+#define MLIR_TARGET_LLVMIR_DIALECT_LLVMARMNEON_LLVMARMNEONTOLLVMIRTRANSLATION_H
+
+#include "mlir/Target/LLVMIR/LLVMTranslationInterface.h"
+
+namespace mlir {
+
+/// Implementation of the dialect interface that converts operations belonging
+/// to the LLVMArmNeon dialect to LLVM IR.
+class LLVMArmNeonDialectLLVMIRTranslationInterface
+ : public LLVMTranslationDialectInterface {
+public:
+ using LLVMTranslationDialectInterface::LLVMTranslationDialectInterface;
+
+ /// Translates the given operation to LLVM IR using the provided IR builder
+ /// and saving the state in `moduleTranslation`.
+ LogicalResult
+ convertOperation(Operation *op, llvm::IRBuilderBase &builder,
+ LLVM::ModuleTranslation &moduleTranslation) const final;
+};
+
+} // namespace mlir
+
+#endif // MLIR_TARGET_LLVMIR_DIALECT_LLVMARMNEON_LLVMARMNEONTOLLVMIRTRANSLATION_H
diff --git a/mlir/include/mlir/Target/LLVMIR/Dialect/LLVMArmSVE/LLVMArmSVEToLLVMIRTranslation.h b/mlir/include/mlir/Target/LLVMIR/Dialect/LLVMArmSVE/LLVMArmSVEToLLVMIRTranslation.h
new file mode 100644
index 000000000000..9d4d05b9b9bd
--- /dev/null
+++ b/mlir/include/mlir/Target/LLVMIR/Dialect/LLVMArmSVE/LLVMArmSVEToLLVMIRTranslation.h
@@ -0,0 +1,37 @@
+//===- LLVMArmSVEToLLVMIRTranslation.h - LLVMArmSVE to LLVM IR --*- C++ -*-===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+//
+// This file implements the dialect interface for translating the LLVMArmSVE
+// dialect to LLVM IR.
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef MLIR_TARGET_LLVMIR_DIALECT_LLVMARMSVE_LLVMARMSVETOLLVMIRTRANSLATION_H
+#define MLIR_TARGET_LLVMIR_DIALECT_LLVMARMSVE_LLVMARMSVETOLLVMIRTRANSLATION_H
+
+#include "mlir/Target/LLVMIR/LLVMTranslationInterface.h"
+
+namespace mlir {
+
+/// Implementation of the dialect interface that converts operations belonging
+/// to the LLVMArmSVE dialect to LLVM IR.
+class LLVMArmSVEDialectLLVMIRTranslationInterface
+ : public LLVMTranslationDialectInterface {
+public:
+ using LLVMTranslationDialectInterface::LLVMTranslationDialectInterface;
+
+ /// Translates the given operation to LLVM IR using the provided IR builder
+ /// and saving the state in `moduleTranslation`.
+ LogicalResult
+ convertOperation(Operation *op, llvm::IRBuilderBase &builder,
+ LLVM::ModuleTranslation &moduleTranslation) const final;
+};
+
+} // namespace mlir
+
+#endif // MLIR_TARGET_LLVMIR_DIALECT_LLVMARMSVE_LLVMARMSVETOLLVMIRTRANSLATION_H
diff --git a/mlir/include/mlir/Target/LLVMIR/Dialect/LLVMIR/LLVMToLLVMIRTranslation.h b/mlir/include/mlir/Target/LLVMIR/Dialect/LLVMIR/LLVMToLLVMIRTranslation.h
index 8b72cedf1ff2..2af76e092917 100644
--- a/mlir/include/mlir/Target/LLVMIR/Dialect/LLVMIR/LLVMToLLVMIRTranslation.h
+++ b/mlir/include/mlir/Target/LLVMIR/Dialect/LLVMIR/LLVMToLLVMIRTranslation.h
@@ -18,8 +18,8 @@
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:
diff --git a/mlir/include/mlir/Target/LLVMIR/Dialect/NVVM/NVVMToLLVMIRTranslation.h b/mlir/include/mlir/Target/LLVMIR/Dialect/NVVM/NVVMToLLVMIRTranslation.h
new file mode 100644
index 000000000000..3a8a01df84b0
--- /dev/null
+++ b/mlir/include/mlir/Target/LLVMIR/Dialect/NVVM/NVVMToLLVMIRTranslation.h
@@ -0,0 +1,42 @@
+//===- NVVMToLLVMIRTranslation.h - NVVM to LLVM IR --------------*- C++ -*-===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+//
+// This file implements the dialect interface for translating the NVVM
+// dialect to LLVM IR.
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef MLIR_TARGET_LLVMIR_DIALECT_NVVM_NVVMTOLLVMIRTRANSLATION_H
+#define MLIR_TARGET_LLVMIR_DIALECT_NVVM_NVVMTOLLVMIRTRANSLATION_H
+
+#include "mlir/Target/LLVMIR/LLVMTranslationInterface.h"
+
+namespace mlir {
+
+/// Implementation of the dialect interface that converts operations belonging
+/// to the NVVM dialect to LLVM IR.
+class NVVMDialectLLVMIRTranslationInterface
+ : public LLVMTranslationDialectInterface {
+public:
+ using LLVMTranslationDialectInterface::LLVMTranslationDialectInterface;
+
+ /// Translates the given operation to LLVM IR using the provided IR builder
+ /// and saving the state in `moduleTranslation`.
+ LogicalResult
+ convertOperation(Operation *op, llvm::IRBuilderBase &builder,
+ LLVM::ModuleTranslation &moduleTranslation) const final;
+
+ /// Attaches module-level metadata for functions marked as kernels.
+ LogicalResult
+ amendOperation(Operation *op, NamedAttribute attribute,
+ LLVM::ModuleTranslation &moduleTranslation) const final;
+};
+
+} // namespace mlir
+
+#endif // MLIR_TARGET_LLVMIR_DIALECT_NVVM_NVVMTOLLVMIRTRANSLATION_H
diff --git a/mlir/include/mlir/Target/LLVMIR/Dialect/OpenMP/OpenMPToLLVMIRTranslation.h b/mlir/include/mlir/Target/LLVMIR/Dialect/OpenMP/OpenMPToLLVMIRTranslation.h
index 7d9eeea9462e..07721d089689 100644
--- a/mlir/include/mlir/Target/LLVMIR/Dialect/OpenMP/OpenMPToLLVMIRTranslation.h
+++ b/mlir/include/mlir/Target/LLVMIR/Dialect/OpenMP/OpenMPToLLVMIRTranslation.h
@@ -18,8 +18,8 @@
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:
diff --git a/mlir/include/mlir/Target/LLVMIR/Dialect/ROCDL/ROCDLToLLVMIRTranslation.h b/mlir/include/mlir/Target/LLVMIR/Dialect/ROCDL/ROCDLToLLVMIRTranslation.h
new file mode 100644
index 000000000000..e2211a59098f
--- /dev/null
+++ b/mlir/include/mlir/Target/LLVMIR/Dialect/ROCDL/ROCDLToLLVMIRTranslation.h
@@ -0,0 +1,42 @@
+//===- ROCDLToLLVMIRTranslation.h - ROCDL to LLVM IR ------------*- C++ -*-===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+//
+// This file implements the dialect interface for translating the ROCDL
+// dialect to LLVM IR.
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef MLIR_TARGET_LLVMIR_DIALECT_ROCDL_ROCDLTOLLVMIRTRANSLATION_H
+#define MLIR_TARGET_LLVMIR_DIALECT_ROCDL_ROCDLTOLLVMIRTRANSLATION_H
+
+#include "mlir/Target/LLVMIR/LLVMTranslationInterface.h"
+
+namespace mlir {
+
+/// Implementation of the dialect interface that converts operations belonging
+/// to the ROCDL dialect to LLVM IR.
+class ROCDLDialectLLVMIRTranslationInterface
+ : public LLVMTranslationDialectInterface {
+public:
+ using LLVMTranslationDialectInterface::LLVMTranslationDialectInterface;
+
+ /// Translates the given operation to LLVM IR using the provided IR builder
+ /// and saving the state in `moduleTranslation`.
+ LogicalResult
+ convertOperation(Operation *op, llvm::IRBuilderBase &builder,
+ LLVM::ModuleTranslation &moduleTranslation) const final;
+
+ /// Attaches module-level metadata for functions marked as kernels.
+ LogicalResult
+ amendOperation(Operation *op, NamedAttribute attribute,
+ LLVM::ModuleTranslation &moduleTranslation) const final;
+};
+
+} // namespace mlir
+
+#endif // MLIR_TARGET_LLVMIR_DIALECT_ROCDL_ROCDLTOLLVMIRTRANSLATION_H
diff --git a/mlir/include/mlir/Target/LLVMIR/LLVMTranslationInterface.h b/mlir/include/mlir/Target/LLVMIR/LLVMTranslationInterface.h
index 0063beac2977..0c563e6e7d39 100644
--- a/mlir/include/mlir/Target/LLVMIR/LLVMTranslationInterface.h
+++ b/mlir/include/mlir/Target/LLVMIR/LLVMTranslationInterface.h
@@ -13,12 +13,14 @@
#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 {
@@ -43,6 +45,18 @@ class LLVMTranslationDialectInterface
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
+ ///
diff erent 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
@@ -61,6 +75,18 @@ class LLVMTranslationInterface
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
diff --git a/mlir/include/mlir/Target/LLVMIR/ModuleTranslation.h b/mlir/include/mlir/Target/LLVMIR/ModuleTranslation.h
index 03b7f5336461..004524f33fa4 100644
--- a/mlir/include/mlir/Target/LLVMIR/ModuleTranslation.h
+++ b/mlir/include/mlir/Target/LLVMIR/ModuleTranslation.h
@@ -142,18 +142,11 @@ class ModuleTranslation {
/// 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.
@@ -184,6 +177,10 @@ class ModuleTranslation {
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
@@ -208,6 +205,9 @@ class ModuleTranslation {
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;
@@ -228,6 +228,8 @@ class ModuleTranslation {
/// 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.
@@ -249,6 +251,20 @@ void connectPHINodes(Region ®ion, const ModuleTranslation &state);
/// Get a topologically sorted list of blocks of the given region.
llvm::SetVector<Block *> getTopologicallySortedBlocks(Region ®ion);
+
+/// 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
diff --git a/mlir/include/mlir/Target/NVVMIR.h b/mlir/include/mlir/Target/NVVMIR.h
deleted file mode 100644
index 0cd7688e275b..000000000000
--- a/mlir/include/mlir/Target/NVVMIR.h
+++ /dev/null
@@ -1,39 +0,0 @@
-//===- 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
diff --git a/mlir/include/mlir/Target/ROCDLIR.h b/mlir/include/mlir/Target/ROCDLIR.h
deleted file mode 100644
index e2cb812a173d..000000000000
--- a/mlir/include/mlir/Target/ROCDLIR.h
+++ /dev/null
@@ -1,40 +0,0 @@
-//===- 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
diff --git a/mlir/lib/Target/CMakeLists.txt b/mlir/lib/Target/CMakeLists.txt
index e951ffade6aa..a23222d37ede 100644
--- a/mlir/lib/Target/CMakeLists.txt
+++ b/mlir/lib/Target/CMakeLists.txt
@@ -24,26 +24,6 @@ add_mlir_translation_library(MLIRTargetLLVMIRModuleTranslation
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
@@ -56,89 +36,12 @@ add_mlir_translation_library(MLIRTargetLLVMIR
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
)
diff --git a/mlir/lib/Target/LLVMIR/ConvertToLLVMIR.cpp b/mlir/lib/Target/LLVMIR/ConvertToLLVMIR.cpp
index 6b30748cc79b..42391513bacf 100644
--- a/mlir/lib/Target/LLVMIR/ConvertToLLVMIR.cpp
+++ b/mlir/lib/Target/LLVMIR/ConvertToLLVMIR.cpp
@@ -12,9 +12,19 @@
#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"
@@ -26,14 +36,14 @@
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;
}
@@ -70,9 +80,24 @@ void registerToLLVMIRTranslation() {
return success();
},
[](DialectRegistry ®istry) {
- 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);
});
}
diff --git a/mlir/lib/Target/LLVMIR/ConvertToNVVMIR.cpp b/mlir/lib/Target/LLVMIR/ConvertToNVVMIR.cpp
deleted file mode 100644
index 7aee913a27d7..000000000000
--- a/mlir/lib/Target/LLVMIR/ConvertToNVVMIR.cpp
+++ /dev/null
@@ -1,124 +0,0 @@
-//===- 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 ®istry) {
- registry.insert<NVVM::NVVMDialect>();
- registerLLVMDialectTranslation(registry);
- });
-}
-} // namespace mlir
diff --git a/mlir/lib/Target/LLVMIR/ConvertToROCDLIR.cpp b/mlir/lib/Target/LLVMIR/ConvertToROCDLIR.cpp
deleted file mode 100644
index 7ebbd3f9ea01..000000000000
--- a/mlir/lib/Target/LLVMIR/ConvertToROCDLIR.cpp
+++ /dev/null
@@ -1,127 +0,0 @@
-//===- 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 ®istry) {
- registry.insert<ROCDL::ROCDLDialect>();
- registerLLVMDialectTranslation(registry);
- });
-}
-} // namespace mlir
diff --git a/mlir/lib/Target/LLVMIR/Dialect/CMakeLists.txt b/mlir/lib/Target/LLVMIR/Dialect/CMakeLists.txt
index 9ab260874b7a..36a28b87403b 100644
--- a/mlir/lib/Target/LLVMIR/Dialect/CMakeLists.txt
+++ b/mlir/lib/Target/LLVMIR/Dialect/CMakeLists.txt
@@ -1,2 +1,7 @@
+add_subdirectory(LLVMArmNeon)
+add_subdirectory(LLVMArmSVE)
+add_subdirectory(LLVMAVX512)
add_subdirectory(LLVMIR)
+add_subdirectory(NVVM)
add_subdirectory(OpenMP)
+add_subdirectory(ROCDL)
diff --git a/mlir/lib/Target/LLVMIR/Dialect/LLVMAVX512/CMakeLists.txt b/mlir/lib/Target/LLVMIR/Dialect/LLVMAVX512/CMakeLists.txt
new file mode 100644
index 000000000000..ba7f594cd81a
--- /dev/null
+++ b/mlir/lib/Target/LLVMIR/Dialect/LLVMAVX512/CMakeLists.txt
@@ -0,0 +1,16 @@
+add_mlir_translation_library(MLIRLLVMAVX512ToLLVMIRTranslation
+ LLVMAVX512ToLLVMIRTranslation.cpp
+
+ DEPENDS
+ MLIRLLVMAVX512ConversionsIncGen
+
+ LINK_COMPONENTS
+ Core
+
+ LINK_LIBS PUBLIC
+ MLIRIR
+ MLIRLLVMAVX512
+ MLIRLLVMIR
+ MLIRSupport
+ MLIRTargetLLVMIRModuleTranslation
+ )
diff --git a/mlir/lib/Target/LLVMIR/Dialect/LLVMAVX512/LLVMAVX512ToLLVMIRTranslation.cpp b/mlir/lib/Target/LLVMIR/Dialect/LLVMAVX512/LLVMAVX512ToLLVMIRTranslation.cpp
new file mode 100644
index 000000000000..dc4742d8e902
--- /dev/null
+++ b/mlir/lib/Target/LLVMIR/Dialect/LLVMAVX512/LLVMAVX512ToLLVMIRTranslation.cpp
@@ -0,0 +1,33 @@
+//===- 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();
+}
diff --git a/mlir/lib/Target/LLVMIR/Dialect/LLVMArmNeon/CMakeLists.txt b/mlir/lib/Target/LLVMIR/Dialect/LLVMArmNeon/CMakeLists.txt
new file mode 100644
index 000000000000..4ff759ad7521
--- /dev/null
+++ b/mlir/lib/Target/LLVMIR/Dialect/LLVMArmNeon/CMakeLists.txt
@@ -0,0 +1,16 @@
+add_mlir_translation_library(MLIRLLVMArmNeonToLLVMIRTranslation
+ LLVMArmNeonToLLVMIRTranslation.cpp
+
+ DEPENDS
+ MLIRLLVMArmNeonConversionsIncGen
+
+ LINK_COMPONENTS
+ Core
+
+ LINK_LIBS PUBLIC
+ MLIRIR
+ MLIRLLVMArmNeon
+ MLIRLLVMIR
+ MLIRSupport
+ MLIRTargetLLVMIRModuleTranslation
+ )
diff --git a/mlir/lib/Target/LLVMIR/Dialect/LLVMArmNeon/LLVMArmNeonToLLVMIRTranslation.cpp b/mlir/lib/Target/LLVMIR/Dialect/LLVMArmNeon/LLVMArmNeonToLLVMIRTranslation.cpp
new file mode 100644
index 000000000000..80776ef19cfc
--- /dev/null
+++ b/mlir/lib/Target/LLVMIR/Dialect/LLVMArmNeon/LLVMArmNeonToLLVMIRTranslation.cpp
@@ -0,0 +1,33 @@
+//===- 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();
+}
diff --git a/mlir/lib/Target/LLVMIR/Dialect/LLVMArmSVE/CMakeLists.txt b/mlir/lib/Target/LLVMIR/Dialect/LLVMArmSVE/CMakeLists.txt
new file mode 100644
index 000000000000..b62391618ba0
--- /dev/null
+++ b/mlir/lib/Target/LLVMIR/Dialect/LLVMArmSVE/CMakeLists.txt
@@ -0,0 +1,16 @@
+add_mlir_translation_library(MLIRLLVMArmSVEToLLVMIRTranslation
+ LLVMArmSVEToLLVMIRTranslation.cpp
+
+ DEPENDS
+ MLIRLLVMArmSVEConversionsIncGen
+
+ LINK_COMPONENTS
+ Core
+
+ LINK_LIBS PUBLIC
+ MLIRIR
+ MLIRLLVMArmSVE
+ MLIRLLVMIR
+ MLIRSupport
+ MLIRTargetLLVMIRModuleTranslation
+ )
diff --git a/mlir/lib/Target/LLVMIR/Dialect/LLVMArmSVE/LLVMArmSVEToLLVMIRTranslation.cpp b/mlir/lib/Target/LLVMIR/Dialect/LLVMArmSVE/LLVMArmSVEToLLVMIRTranslation.cpp
new file mode 100644
index 000000000000..7bd9566b0972
--- /dev/null
+++ b/mlir/lib/Target/LLVMIR/Dialect/LLVMArmSVE/LLVMArmSVEToLLVMIRTranslation.cpp
@@ -0,0 +1,33 @@
+//===- 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();
+}
diff --git a/mlir/lib/Target/LLVMIR/Dialect/LLVMIR/LLVMToLLVMIRTranslation.cpp b/mlir/lib/Target/LLVMIR/Dialect/LLVMIR/LLVMToLLVMIRTranslation.cpp
index 25d529439423..c14eba8b6ea0 100644
--- a/mlir/lib/Target/LLVMIR/Dialect/LLVMIR/LLVMToLLVMIRTranslation.cpp
+++ b/mlir/lib/Target/LLVMIR/Dialect/LLVMIR/LLVMToLLVMIRTranslation.cpp
@@ -23,6 +23,7 @@
using namespace mlir;
using namespace mlir::LLVM;
+using mlir::LLVM::detail::getLLVMConstant;
#include "mlir/Dialect/LLVMIR/LLVMConversionEnumsToLLVM.inc"
@@ -166,27 +167,6 @@ static llvm::FastMathFlags getFastmathFlags(FastmathFlagsInterface &op) {
return ret;
}
-namespace {
-/// Dispatcher functional object targeting
diff erent 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) {
@@ -202,21 +182,6 @@ convertOperationImpl(Operation &opInst, llvm::IRBuilderBase &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
@@ -243,7 +208,7 @@ convertOperationImpl(Operation &opInst, llvm::IRBuilderBase &builder,
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.
@@ -269,23 +234,25 @@ convertOperationImpl(Operation &opInst, llvm::IRBuilderBase &builder,
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()),
@@ -306,17 +273,18 @@ convertOperationImpl(Operation &opInst, llvm::IRBuilderBase &builder,
}
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();
}
@@ -365,8 +333,8 @@ convertOperationImpl(Operation &opInst, llvm::IRBuilderBase &builder,
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()))
@@ -389,9 +357,10 @@ convertOperationImpl(Operation &opInst, llvm::IRBuilderBase &builder,
assert((global || function) &&
"referencing an undefined global or function");
- mapValue(addressOfOp.getResult(),
- global ? moduleTranslation.lookupGlobal(global)
- : moduleTranslation.lookupFunction(function.getName()));
+ moduleTranslation.mapValue(
+ addressOfOp.getResult(),
+ global ? moduleTranslation.lookupGlobal(global)
+ : moduleTranslation.lookupFunction(function.getName()));
return success();
}
diff --git a/mlir/lib/Target/LLVMIR/Dialect/NVVM/CMakeLists.txt b/mlir/lib/Target/LLVMIR/Dialect/NVVM/CMakeLists.txt
new file mode 100644
index 000000000000..5a2615a84a38
--- /dev/null
+++ b/mlir/lib/Target/LLVMIR/Dialect/NVVM/CMakeLists.txt
@@ -0,0 +1,16 @@
+add_mlir_translation_library(MLIRNVVMToLLVMIRTranslation
+ NVVMToLLVMIRTranslation.cpp
+
+ DEPENDS
+ MLIRNVVMConversionsIncGen
+
+ LINK_COMPONENTS
+ Core
+
+ LINK_LIBS PUBLIC
+ MLIRIR
+ MLIRLLVMIR
+ MLIRNVVMIR
+ MLIRSupport
+ MLIRTargetLLVMIRModuleTranslation
+ )
diff --git a/mlir/lib/Target/LLVMIR/Dialect/NVVM/NVVMToLLVMIRTranslation.cpp b/mlir/lib/Target/LLVMIR/Dialect/NVVM/NVVMToLLVMIRTranslation.cpp
new file mode 100644
index 000000000000..fdc9add52dee
--- /dev/null
+++ b/mlir/lib/Target/LLVMIR/Dialect/NVVM/NVVMToLLVMIRTranslation.cpp
@@ -0,0 +1,67 @@
+//===- 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();
+}
diff --git a/mlir/lib/Target/LLVMIR/Dialect/ROCDL/CMakeLists.txt b/mlir/lib/Target/LLVMIR/Dialect/ROCDL/CMakeLists.txt
new file mode 100644
index 000000000000..7e542ec77adb
--- /dev/null
+++ b/mlir/lib/Target/LLVMIR/Dialect/ROCDL/CMakeLists.txt
@@ -0,0 +1,16 @@
+add_mlir_translation_library(MLIRROCDLToLLVMIRTranslation
+ ROCDLToLLVMIRTranslation.cpp
+
+ DEPENDS
+ MLIRROCDLConversionsIncGen
+
+ LINK_COMPONENTS
+ Core
+
+ LINK_LIBS PUBLIC
+ MLIRIR
+ MLIRLLVMIR
+ MLIRROCDLIR
+ MLIRSupport
+ MLIRTargetLLVMIRModuleTranslation
+ )
diff --git a/mlir/lib/Target/LLVMIR/Dialect/ROCDL/ROCDLToLLVMIRTranslation.cpp b/mlir/lib/Target/LLVMIR/Dialect/ROCDL/ROCDLToLLVMIRTranslation.cpp
new file mode 100644
index 000000000000..7b34f803e7e3
--- /dev/null
+++ b/mlir/lib/Target/LLVMIR/Dialect/ROCDL/ROCDLToLLVMIRTranslation.cpp
@@ -0,0 +1,69 @@
+//===- 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();
+}
diff --git a/mlir/lib/Target/LLVMIR/LLVMAVX512Intr.cpp b/mlir/lib/Target/LLVMIR/LLVMAVX512Intr.cpp
deleted file mode 100644
index 082504d65779..000000000000
--- a/mlir/lib/Target/LLVMIR/LLVMAVX512Intr.cpp
+++ /dev/null
@@ -1,65 +0,0 @@
-//===- 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 ®istry) {
- registry.insert<LLVM::LLVMAVX512Dialect>();
- registerLLVMDialectTranslation(registry);
- });
-}
-} // namespace mlir
diff --git a/mlir/lib/Target/LLVMIR/LLVMArmNeonIntr.cpp b/mlir/lib/Target/LLVMIR/LLVMArmNeonIntr.cpp
deleted file mode 100644
index 14bbd14908d1..000000000000
--- a/mlir/lib/Target/LLVMIR/LLVMArmNeonIntr.cpp
+++ /dev/null
@@ -1,65 +0,0 @@
-//===- 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 ®istry) {
- registry.insert<LLVM::LLVMArmNeonDialect>();
- registerLLVMDialectTranslation(registry);
- });
-}
-} // namespace mlir
diff --git a/mlir/lib/Target/LLVMIR/LLVMArmSVEIntr.cpp b/mlir/lib/Target/LLVMIR/LLVMArmSVEIntr.cpp
deleted file mode 100644
index 5ef8b48c709e..000000000000
--- a/mlir/lib/Target/LLVMIR/LLVMArmSVEIntr.cpp
+++ /dev/null
@@ -1,65 +0,0 @@
-//===- 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 ®istry) {
- registry.insert<LLVM::LLVMArmSVEDialect>();
- registerLLVMDialectTranslation(registry);
- });
-}
-} // namespace mlir
diff --git a/mlir/lib/Target/LLVMIR/ModuleTranslation.cpp b/mlir/lib/Target/LLVMIR/ModuleTranslation.cpp
index f563c67270e6..84f70a8cc51e 100644
--- a/mlir/lib/Target/LLVMIR/ModuleTranslation.cpp
+++ b/mlir/lib/Target/LLVMIR/ModuleTranslation.cpp
@@ -102,9 +102,9 @@ static llvm::Type *getInnermostElementType(llvm::Type *type) {
/// 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()) {
@@ -120,8 +120,8 @@ llvm::Constant *ModuleTranslation::getLLVMConstant(llvm::Type *llvmType,
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;
@@ -140,7 +140,8 @@ llvm::Constant *ModuleTranslation::getLLVMConstant(llvm::Type *llvmType,
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())
@@ -164,7 +165,8 @@ llvm::Constant *ModuleTranslation::getLLVMConstant(llvm::Type *llvmType,
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;
}
@@ -177,8 +179,9 @@ llvm::Constant *ModuleTranslation::getLLVMConstant(llvm::Type *llvmType,
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(),
+ stringAttr.getValue().size()});
}
emitError(loc, "unsupported constant value");
return nullptr;
@@ -288,17 +291,23 @@ mlir::LLVM::detail::getTopologicallySortedBlocks(Region ®ion) {
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
diff erent functions with a
-///
diff erent 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 opInst.emitError("unsupported or non-LLVM operation: ")
- << opInst.getName();
+ return convertDialectAttributes(&opInst);
}
/// Convert block to LLVM IR. Unless `ignoreArguments` is set, emit PHI nodes
@@ -360,8 +369,8 @@ LogicalResult ModuleTranslation::convertGlobals() {
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()) {
@@ -537,7 +546,7 @@ LogicalResult ModuleTranslation::convertOneFunction(LLVMFuncOp func) {
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);
}
@@ -558,9 +567,18 @@ LogicalResult ModuleTranslation::convertOneFunction(LLVMFuncOp func) {
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();
}
@@ -625,6 +643,11 @@ 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>();
diff --git a/mlir/test/Target/arm-neon.mlir b/mlir/test/Target/arm-neon.mlir
index 8e24b5bad295..04627a087d35 100644
--- a/mlir/test/Target/arm-neon.mlir
+++ b/mlir/test/Target/arm-neon.mlir
@@ -1,4 +1,4 @@
-// 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>)> {
diff --git a/mlir/test/Target/arm-sve.mlir b/mlir/test/Target/arm-sve.mlir
index f00992e05bfd..a274071abd36 100644
--- a/mlir/test/Target/arm-sve.mlir
+++ b/mlir/test/Target/arm-sve.mlir
@@ -1,4 +1,4 @@
-// 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>,
diff --git a/mlir/test/Target/avx512.mlir b/mlir/test/Target/avx512.mlir
index 037f44f0449c..69c492766f9d 100644
--- a/mlir/test/Target/avx512.mlir
+++ b/mlir/test/Target/avx512.mlir
@@ -1,4 +1,4 @@
-// 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>,
diff --git a/mlir/test/Target/nvvmir.mlir b/mlir/test/Target/nvvmir.mlir
index 3065188925f8..5dcdeb880de3 100644
--- a/mlir/test/Target/nvvmir.mlir
+++ b/mlir/test/Target/nvvmir.mlir
@@ -1,4 +1,4 @@
-// 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()
diff --git a/mlir/test/Target/rocdl.mlir b/mlir/test/Target/rocdl.mlir
index 27884c738539..38ef3d45a022 100644
--- a/mlir/test/Target/rocdl.mlir
+++ b/mlir/test/Target/rocdl.mlir
@@ -1,4 +1,4 @@
-// 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
diff --git a/mlir/test/lib/Transforms/CMakeLists.txt b/mlir/test/lib/Transforms/CMakeLists.txt
index 124edb08d672..4917616bbd4a 100644
--- a/mlir/test/lib/Transforms/CMakeLists.txt
+++ b/mlir/test/lib/Transforms/CMakeLists.txt
@@ -46,23 +46,24 @@ add_mlir_library(MLIRTestTransforms
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)
diff --git a/mlir/test/lib/Transforms/TestConvertGPUKernelToCubin.cpp b/mlir/test/lib/Transforms/TestConvertGPUKernelToCubin.cpp
index ac217a2956d7..d7da76ba998d 100644
--- a/mlir/test/lib/Transforms/TestConvertGPUKernelToCubin.cpp
+++ b/mlir/test/lib/Transforms/TestConvertGPUKernelToCubin.cpp
@@ -7,9 +7,11 @@
//===----------------------------------------------------------------------===//
#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;
@@ -21,6 +23,27 @@ 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) {
+ registerLLVMDialectTranslation(*m->getContext());
+ registerNVVMDialectTranslation(*m->getContext());
+ return translateModuleToLLVMIR(m, llvmContext, moduleName);
+}
+
namespace mlir {
namespace test {
void registerTestConvertGPUKernelToCubinPass() {
diff --git a/mlir/test/lib/Transforms/TestConvertGPUKernelToHsaco.cpp b/mlir/test/lib/Transforms/TestConvertGPUKernelToHsaco.cpp
index f7036c92ec05..1e559a74cf55 100644
--- a/mlir/test/lib/Transforms/TestConvertGPUKernelToHsaco.cpp
+++ b/mlir/test/lib/Transforms/TestConvertGPUKernelToHsaco.cpp
@@ -7,9 +7,11 @@
//===----------------------------------------------------------------------===//
#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;
@@ -21,6 +23,22 @@ 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) {
+ registerLLVMDialectTranslation(*m->getContext());
+ registerROCDLDialectTranslation(*m->getContext());
+ return translateModuleToLLVMIR(m, llvmContext, moduleName);
+}
+
namespace mlir {
namespace test {
void registerTestConvertGPUKernelToHsacoPass() {
@@ -35,7 +53,7 @@ void registerTestConvertGPUKernelToHsacoPass() {
LLVMInitializeAMDGPUAsmPrinter();
pm.addPass(createConvertGPUKernelToBlobPass(
- translateModuleToROCDLIR, compileIsaToHsacoForTesting,
+ translateModuleToROCDL, compileIsaToHsacoForTesting,
"amdgcn-amd-amdhsa", "gfx900", "-code-object-v3", "rocdl.hsaco"));
});
}
diff --git a/mlir/tools/mlir-cuda-runner/mlir-cuda-runner.cpp b/mlir/tools/mlir-cuda-runner/mlir-cuda-runner.cpp
index 1e6a80b20fa6..4800c2c2fd81 100644
--- a/mlir/tools/mlir-cuda-runner/mlir-cuda-runner.cpp
+++ b/mlir/tools/mlir-cuda-runner/mlir-cuda-runner.cpp
@@ -32,7 +32,7 @@
#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"
@@ -121,7 +121,7 @@ static LogicalResult runMLIRPasses(ModuleOp m) {
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());
@@ -156,6 +156,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);
return mlir::JitRunnerMain(argc, argv, registry, jitRunnerConfig);
diff --git a/mlir/tools/mlir-rocm-runner/mlir-rocm-runner.cpp b/mlir/tools/mlir-rocm-runner/mlir-rocm-runner.cpp
index 8237545cc00b..14251b4c3874 100644
--- a/mlir/tools/mlir-rocm-runner/mlir-rocm-runner.cpp
+++ b/mlir/tools/mlir-rocm-runner/mlir-rocm-runner.cpp
@@ -31,6 +31,7 @@
#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"
@@ -341,6 +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);
return mlir::JitRunnerMain(argc, argv, registry, jitRunnerConfig);
diff --git a/mlir/tools/mlir-tblgen/LLVMIRConversionGen.cpp b/mlir/tools/mlir-tblgen/LLVMIRConversionGen.cpp
index 8384bbffe26c..42b053a9ed64 100644
--- a/mlir/tools/mlir-tblgen/LLVMIRConversionGen.cpp
+++ b/mlir/tools/mlir-tblgen/LLVMIRConversionGen.cpp
@@ -126,15 +126,17 @@ static bool emitOneBuilder(const Record &record, raw_ostream &os) {
// 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") {
More information about the Mlir-commits
mailing list