maxjhandsome
/
llvm-project
forked from OSchip/llvm-project
1
0
Fork 0
Code Issues Pull Requests Packages Releases Wiki Activity

[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
This commit is contained in:
Alex Zinenko 2021-02-12 12:53:27 +01:00
parent 2d728bbff5
commit 176379e0c8
48 changed files with 750 additions and 760 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

3
mlir/examples/standalone/test/Standalone/standalone-translate.mlir
View File

@ -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

13
mlir/include/mlir/Dialect/LLVMIR/LLVMOpBase.td
View File

@ -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);

3
mlir/include/mlir/Dialect/LLVMIR/LLVMOps.td
View File

@ -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 = [{

2
mlir/include/mlir/Dialect/LLVMIR/ROCDLOps.td
View File

@ -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});

10
mlir/include/mlir/InitAllTranslations.h
View File

@ -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;

4
mlir/include/mlir/Target/LLVMIR.h
View File

@ -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

37
mlir/include/mlir/Target/LLVMIR/Dialect/LLVMAVX512/LLVMAVX512ToLLVMIRTranslation.h Normal file
View File

@ -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

37
mlir/include/mlir/Target/LLVMIR/Dialect/LLVMArmNeon/LLVMArmNeonToLLVMIRTranslation.h Normal file
View File

@ -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

37
mlir/include/mlir/Target/LLVMIR/Dialect/LLVMArmSVE/LLVMArmSVEToLLVMIRTranslation.h Normal file
View File

@ -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

4
mlir/include/mlir/Target/LLVMIR/Dialect/LLVMIR/LLVMToLLVMIRTranslation.h
View File

@ -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:

42
mlir/include/mlir/Target/LLVMIR/Dialect/NVVM/NVVMToLLVMIRTranslation.h Normal file
View File

@ -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

4
mlir/include/mlir/Target/LLVMIR/Dialect/OpenMP/OpenMPToLLVMIRTranslation.h
View File

@ -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:

42
mlir/include/mlir/Target/LLVMIR/Dialect/ROCDL/ROCDLToLLVMIRTranslation.h Normal file
View File

@ -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

28
mlir/include/mlir/Target/LLVMIR/LLVMTranslationInterface.h
View File

@ -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 @@ public:
LLVM::ModuleTranslation &moduleTranslation) const {
return failure();
}
/// Hook for derived dialect interface to act on an operation that has dialect
/// attributes from the derived dialect (the operation itself may be from a
/// different dialect). This gets called after the operation has been
/// translated. The hook is expected to use moduleTranslation to look up the
/// translation results and amend the corresponding IR constructs. Does
/// nothing and succeeds by default.
virtual LogicalResult
amendOperation(Operation *op, NamedAttribute attribute,
LLVM::ModuleTranslation &moduleTranslation) const {
return success();
}
};
/// Interface collection for translation to LLVM IR, dispatches to a concrete
@ -61,6 +75,18 @@ public:
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

32
mlir/include/mlir/Target/LLVMIR/ModuleTranslation.h
View File

@ -142,18 +142,11 @@ public:
/// 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 @@ public:
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 @@ private:
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 @@ private:
/// 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 &region, const ModuleTranslation &state);
/// Get a topologically sorted list of blocks of the given region.
llvm::SetVector<Block *> getTopologicallySortedBlocks(Region &region);
/// Create an LLVM IR constant of `llvmType` from the MLIR attribute `attr`.
/// This currently supports integer, floating point, splat and dense element
/// attributes and combinations thereof. In case of error, report it to `loc`
/// and return nullptr.
llvm::Constant *getLLVMConstant(llvm::Type *llvmType, Attribute attr,
Location loc,
const ModuleTranslation &moduleTranslation);
/// Creates a call to an LLVM IR intrinsic function with the given arguments.
llvm::Value *createIntrinsicCall(llvm::IRBuilderBase &builder,
llvm::Intrinsic::ID intrinsic,
ArrayRef<llvm::Value *> args = {},
ArrayRef<llvm::Type *> tys = {});
} // namespace detail
} // namespace LLVM

39
mlir/include/mlir/Target/NVVMIR.h
View File

@ -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

40
mlir/include/mlir/Target/ROCDLIR.h
View File

@ -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

107
mlir/lib/Target/CMakeLists.txt
View File

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

35
mlir/lib/Target/LLVMIR/ConvertToLLVMIR.cpp
View File

@ -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 &registry) {
registry.insert<omp::OpenMPDialect>();
registry.insert<omp::OpenMPDialect, LLVM::LLVMAVX512Dialect,
LLVM::LLVMArmSVEDialect, LLVM::LLVMArmNeonDialect,
NVVM::NVVMDialect, ROCDL::ROCDLDialect>();
registry.addDialectInterface<omp::OpenMPDialect,
OpenMPDialectLLVMIRTranslationInterface>();
registry
.addDialectInterface<LLVM::LLVMAVX512Dialect,
LLVMAVX512DialectLLVMIRTranslationInterface>();
registry.addDialectInterface<
LLVM::LLVMArmNeonDialect,
LLVMArmNeonDialectLLVMIRTranslationInterface>();
registry
.addDialectInterface<LLVM::LLVMArmSVEDialect,
LLVMArmSVEDialectLLVMIRTranslationInterface>();
registry.addDialectInterface<NVVM::NVVMDialect,
NVVMDialectLLVMIRTranslationInterface>();
registry.addDialectInterface<ROCDL::ROCDLDialect,
ROCDLDialectLLVMIRTranslationInterface>();
registerLLVMDialectTranslation(registry);
});
}

124
mlir/lib/Target/LLVMIR/ConvertToNVVMIR.cpp
View File

@ -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 &registry) {
registry.insert<NVVM::NVVMDialect>();
registerLLVMDialectTranslation(registry);
});
}
} // namespace mlir

127
mlir/lib/Target/LLVMIR/ConvertToROCDLIR.cpp
View File

@ -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 &registry) {
registry.insert<ROCDL::ROCDLDialect>();
registerLLVMDialectTranslation(registry);
});
}
} // namespace mlir

5
mlir/lib/Target/LLVMIR/Dialect/CMakeLists.txt
View File

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

16
mlir/lib/Target/LLVMIR/Dialect/LLVMAVX512/CMakeLists.txt Normal file
View File

@ -0,0 +1,16 @@
add_mlir_translation_library(MLIRLLVMAVX512ToLLVMIRTranslation
LLVMAVX512ToLLVMIRTranslation.cpp
DEPENDS
MLIRLLVMAVX512ConversionsIncGen
LINK_COMPONENTS
Core
LINK_LIBS PUBLIC
MLIRIR
MLIRLLVMAVX512
MLIRLLVMIR
MLIRSupport
MLIRTargetLLVMIRModuleTranslation
)

33
mlir/lib/Target/LLVMIR/Dialect/LLVMAVX512/LLVMAVX512ToLLVMIRTranslation.cpp Normal file
View File

@ -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();
}

16
mlir/lib/Target/LLVMIR/Dialect/LLVMArmNeon/CMakeLists.txt Normal file
View File

@ -0,0 +1,16 @@
add_mlir_translation_library(MLIRLLVMArmNeonToLLVMIRTranslation
LLVMArmNeonToLLVMIRTranslation.cpp
DEPENDS
MLIRLLVMArmNeonConversionsIncGen
LINK_COMPONENTS
Core
LINK_LIBS PUBLIC
MLIRIR
MLIRLLVMArmNeon
MLIRLLVMIR
MLIRSupport
MLIRTargetLLVMIRModuleTranslation
)

33
mlir/lib/Target/LLVMIR/Dialect/LLVMArmNeon/LLVMArmNeonToLLVMIRTranslation.cpp Normal file
View File

@ -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();
}

16
mlir/lib/Target/LLVMIR/Dialect/LLVMArmSVE/CMakeLists.txt Normal file
View File

@ -0,0 +1,16 @@
add_mlir_translation_library(MLIRLLVMArmSVEToLLVMIRTranslation
LLVMArmSVEToLLVMIRTranslation.cpp
DEPENDS
MLIRLLVMArmSVEConversionsIncGen
LINK_COMPONENTS
Core
LINK_LIBS PUBLIC
MLIRIR
MLIRLLVMArmSVE
MLIRLLVMIR
MLIRSupport
MLIRTargetLLVMIRModuleTranslation
)

33
mlir/lib/Target/LLVMIR/Dialect/LLVMArmSVE/LLVMArmSVEToLLVMIRTranslation.cpp Normal file
View File

@ -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();
}

71
mlir/lib/Target/LLVMIR/Dialect/LLVMIR/LLVMToLLVMIRTranslation.cpp
View File

@ -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 different overloads of
/// ModuleTranslation::mapValue.
// TODO: this is only necessary for compatibility with the code emitted from
// ODS, remove when ODS is updated (after all dialects have migrated to the new
// translation mechanism).
struct MapValueDispatcher {
explicit MapValueDispatcher(ModuleTranslation &mt) : moduleTranslation(mt) {}
llvm::Value *&operator()(mlir::Value v) {
return moduleTranslation.mapValue(v);
}
void operator()(mlir::Value m, llvm::Value *l) {
moduleTranslation.mapValue(m, l);
}
LLVM::ModuleTranslation &moduleTranslation;
};
} // end namespace
static LogicalResult
convertOperationImpl(Operation &opInst, llvm::IRBuilderBase &builder,
LLVM::ModuleTranslation &moduleTranslation) {
@ -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();
}

16
mlir/lib/Target/LLVMIR/Dialect/NVVM/CMakeLists.txt Normal file
View File

@ -0,0 +1,16 @@
add_mlir_translation_library(MLIRNVVMToLLVMIRTranslation
NVVMToLLVMIRTranslation.cpp
DEPENDS
MLIRNVVMConversionsIncGen
LINK_COMPONENTS
Core
LINK_LIBS PUBLIC
MLIRIR
MLIRLLVMIR
MLIRNVVMIR
MLIRSupport
MLIRTargetLLVMIRModuleTranslation
)

67
mlir/lib/Target/LLVMIR/Dialect/NVVM/NVVMToLLVMIRTranslation.cpp Normal file
View File

@ -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();
}

16
mlir/lib/Target/LLVMIR/Dialect/ROCDL/CMakeLists.txt Normal file
View File

@ -0,0 +1,16 @@
add_mlir_translation_library(MLIRROCDLToLLVMIRTranslation
ROCDLToLLVMIRTranslation.cpp
DEPENDS
MLIRROCDLConversionsIncGen
LINK_COMPONENTS
Core
LINK_LIBS PUBLIC
MLIRIR
MLIRLLVMIR
MLIRROCDLIR
MLIRSupport
MLIRTargetLLVMIRModuleTranslation
)

69
mlir/lib/Target/LLVMIR/Dialect/ROCDL/ROCDLToLLVMIRTranslation.cpp Normal file
View File

@ -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();
}

65
mlir/lib/Target/LLVMIR/LLVMAVX512Intr.cpp
View File

@ -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 &registry) {
registry.insert<LLVM::LLVMAVX512Dialect>();
registerLLVMDialectTranslation(registry);
});
}
} // namespace mlir

65
mlir/lib/Target/LLVMIR/LLVMArmNeonIntr.cpp
View File

@ -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 &registry) {
registry.insert<LLVM::LLVMArmNeonDialect>();
registerLLVMDialectTranslation(registry);
});
}
} // namespace mlir

65
mlir/lib/Target/LLVMIR/LLVMArmSVEIntr.cpp
View File

@ -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 &registry) {
registry.insert<LLVM::LLVMArmSVEDialect>();
registerLLVMDialectTranslation(registry);
});
}
} // namespace mlir

65
mlir/lib/Target/LLVMIR/ModuleTranslation.cpp
View File

@ -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 &region) {
return blocks;
}
llvm::Value *mlir::LLVM::detail::createIntrinsicCall(
llvm::IRBuilderBase &builder, llvm::Intrinsic::ID intrinsic,
ArrayRef<llvm::Value *> args, ArrayRef<llvm::Type *> tys) {
llvm::Module *module = builder.GetInsertBlock()->getModule();
llvm::Function *fn = llvm::Intrinsic::getDeclaration(module, intrinsic, tys);
return builder.CreateCall(fn, args);
}
/// Given a single MLIR operation, create the corresponding LLVM IR operation
/// using the `builder`. LLVM IR Builder does not have a generic interface so
/// this has to be a long chain of `if`s calling different functions with a
/// different number of arguments.
/// using the `builder`.
LogicalResult ModuleTranslation::convertOperation(Operation &opInst,
llvm::IRBuilder<> &builder) {
if (succeeded(iface.convertOperation(&opInst, builder, *this)))
return success();
if (failed(iface.convertOperation(&opInst, builder, *this)))
return opInst.emitError("unsupported or non-LLVM operation: ")
<< opInst.getName();
return 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>();

2
mlir/test/Target/arm-neon.mlir
View File

@ -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>)> {

2
mlir/test/Target/arm-sve.mlir
View File

@ -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>,

2
mlir/test/Target/avx512.mlir
View File

@ -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>,

2
mlir/test/Target/nvvmir.mlir
View File

@ -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()

2
mlir/test/Target/rocdl.mlir
View File

@ -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

13
mlir/test/lib/Transforms/CMakeLists.txt
View File

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

25
mlir/test/lib/Transforms/TestConvertGPUKernelToCubin.cpp
View File

@ -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() {

22
mlir/test/lib/Transforms/TestConvertGPUKernelToHsaco.cpp
View File

@ -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"));
});
}

6
mlir/tools/mlir-cuda-runner/mlir-cuda-runner.cpp
View File

@ -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);

3
mlir/tools/mlir-rocm-runner/mlir-rocm-runner.cpp
View File

@ -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);

10
mlir/tools/mlir-tblgen/LLVMIRConversionGen.cpp
View File

@ -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") {