[Mlir-commits] [mlir] a825fb2 - [mlir] Remove mlir-rocm-runner
Christian Sigg
llvmlistbot at llvm.org
Fri Mar 19 00:24:20 PDT 2021
Author: Christian Sigg
Date: 2021-03-19T00:24:10-07:00
New Revision: a825fb2c07337cc2c84783558e91416e07adcf42
URL: https://github.com/llvm/llvm-project/commit/a825fb2c07337cc2c84783558e91416e07adcf42
DIFF: https://github.com/llvm/llvm-project/commit/a825fb2c07337cc2c84783558e91416e07adcf42.diff
LOG: [mlir] Remove mlir-rocm-runner
This change combines for ROCm what was done for CUDA in D97463, D98203, D98360, and D98396.
I did not try to compile SerializeToHsaco.cpp or test mlir/test/Integration/GPU/ROCM because I don't have an AMD card. I fixed the things that had obvious bit-rot though.
Reviewed By: whchung
Differential Revision: https://reviews.llvm.org/D98447
Added:
mlir/lib/Dialect/GPU/Transforms/SerializeToHsaco.cpp
mlir/lib/ExecutionEngine/RocmRuntimeWrappers.cpp
mlir/test/Integration/GPU/ROCM/gpu-to-hsaco.mlir
mlir/test/Integration/GPU/ROCM/lit.local.cfg
mlir/test/Integration/GPU/ROCM/two-modules.mlir
mlir/test/Integration/GPU/ROCM/vecadd.mlir
mlir/test/Integration/GPU/ROCM/vector-transferops.mlir
Modified:
mlir/include/mlir/Dialect/GPU/Passes.h
mlir/include/mlir/InitAllPasses.h
mlir/lib/Dialect/GPU/CMakeLists.txt
mlir/lib/ExecutionEngine/CMakeLists.txt
mlir/test/CMakeLists.txt
mlir/test/Conversion/GPUToROCm/lower-rocdl-kernel-to-hsaco.mlir
mlir/test/Integration/GPU/CUDA/lit.local.cfg
mlir/test/lib/Transforms/TestConvertGPUKernelToHsaco.cpp
mlir/test/lit.cfg.py
mlir/test/lit.site.cfg.py.in
mlir/tools/CMakeLists.txt
mlir/tools/mlir-opt/mlir-opt.cpp
Removed:
mlir/test/mlir-rocm-runner/gpu-to-hsaco.mlir
mlir/test/mlir-rocm-runner/lit.local.cfg
mlir/test/mlir-rocm-runner/two-modules.mlir
mlir/test/mlir-rocm-runner/vecadd.mlir
mlir/test/mlir-rocm-runner/vector-transferops.mlir
mlir/tools/mlir-rocm-runner/CMakeLists.txt
mlir/tools/mlir-rocm-runner/mlir-rocm-runner.cpp
mlir/tools/mlir-rocm-runner/rocm-runtime-wrappers.cpp
################################################################################
diff --git a/mlir/include/mlir/Dialect/GPU/Passes.h b/mlir/include/mlir/Dialect/GPU/Passes.h
index 6a6a2c0678b6..bfb5626fca19 100644
--- a/mlir/include/mlir/Dialect/GPU/Passes.h
+++ b/mlir/include/mlir/Dialect/GPU/Passes.h
@@ -90,6 +90,10 @@ class SerializeToBlobPass : public OperationPass<gpu::GPUModuleOp> {
/// annotation.
void registerGpuSerializeToCubinPass();
+/// Register pass to serialize GPU kernel functions to a HSAco binary
+/// annotation.
+void registerGpuSerializeToHsacoPass();
+
/// Generate the code for registering passes.
#define GEN_PASS_REGISTRATION
#include "mlir/Dialect/GPU/Passes.h.inc"
diff --git a/mlir/include/mlir/InitAllPasses.h b/mlir/include/mlir/InitAllPasses.h
index 029df0735959..ab9629ac86c1 100644
--- a/mlir/include/mlir/InitAllPasses.h
+++ b/mlir/include/mlir/InitAllPasses.h
@@ -52,6 +52,7 @@ inline void registerAllPasses() {
registerAsyncPasses();
registerGPUPasses();
registerGpuSerializeToCubinPass();
+ registerGpuSerializeToHsacoPass();
registerLinalgPasses();
LLVM::registerLLVMPasses();
quant::registerQuantPasses();
diff --git a/mlir/lib/Dialect/GPU/CMakeLists.txt b/mlir/lib/Dialect/GPU/CMakeLists.txt
index d7fbfe0b5b61..ea70029c849e 100644
--- a/mlir/lib/Dialect/GPU/CMakeLists.txt
+++ b/mlir/lib/Dialect/GPU/CMakeLists.txt
@@ -6,6 +6,16 @@ if (MLIR_CUDA_CONVERSIONS_ENABLED)
)
endif()
+if (MLIR_ROCM_CONVERSIONS_ENABLED)
+ set(AMDGPU_LIBS
+ MCParser
+ AMDGPUAsmParser
+ AMDGPUCodeGen
+ AMDGPUDesc
+ AMDGPUInfo
+ )
+endif()
+
add_mlir_dialect_library(MLIRGPU
IR/GPUDialect.cpp
Transforms/AllReduceLowering.cpp
@@ -15,6 +25,7 @@ add_mlir_dialect_library(MLIRGPU
Transforms/ParallelLoopMapper.cpp
Transforms/SerializeToBlob.cpp
Transforms/SerializeToCubin.cpp
+ Transforms/SerializeToHsaco.cpp
ADDITIONAL_HEADER_DIRS
${MLIR_MAIN_INCLUDE_DIR}/mlir/Dialect/GPU
@@ -23,6 +34,7 @@ add_mlir_dialect_library(MLIRGPU
Core
MC
${NVPTX_LIBS}
+ ${AMDGPU_LIBS}
DEPENDS
MLIRGPUOpsIncGen
@@ -84,3 +96,58 @@ if(MLIR_CUDA_RUNNER_ENABLED)
)
endif()
+
+if(MLIR_ROCM_RUNNER_ENABLED)
+ if (NOT ("AMDGPU" IN_LIST LLVM_TARGETS_TO_BUILD))
+ message(SEND_ERROR
+ "Building mlir with ROCm support requires the AMDGPU backend")
+ endif()
+
+ # Ensure lld is enabled.
+ if (NOT "lld" IN_LIST LLVM_ENABLE_PROJECTS)
+ message(SEND_ERROR "lld is not enabled. Please revise LLVM_ENABLE_PROJECTS")
+ endif()
+
+ # Configure ROCm support.
+ if (NOT DEFINED ROCM_PATH)
+ if (NOT DEFINED ENV{ROCM_PATH})
+ set(ROCM_PATH "/opt/rocm" CACHE PATH "Path to which ROCm has been installed")
+ else()
+ set(ROCM_PATH $ENV{ROCM_PATH} CACHE PATH "Path to which ROCm has been installed")
+ endif()
+ set(HIP_PATH "${ROCM_PATH}/hip" CACHE PATH " Path to which HIP has been installed")
+ endif()
+ set(CMAKE_MODULE_PATH "${HIP_PATH}/cmake" ${CMAKE_MODULE_PATH})
+ find_package(HIP)
+ if (NOT HIP_FOUND)
+ message(SEND_ERROR "Building mlir with ROCm support requires a working ROCm and HIP install")
+ else()
+ message(STATUS "ROCm HIP version: ${HIP_VERSION}")
+ endif()
+
+ target_compile_definitions(obj.MLIRGPU
+ PRIVATE
+ __HIP_PLATFORM_HCC__
+ __ROCM_PATH__="${ROCM_PATH}"
+ MLIR_GPU_TO_HSACO_PASS_ENABLE=1
+ )
+
+ target_include_directories(obj.MLIRGPU
+ PRIVATE
+ ${MLIR_SOURCE_DIR}/../lld/include
+ ${HIP_PATH}/include
+ ${ROCM_PATH}/include
+ )
+
+ target_link_libraries(MLIRGPU
+ PRIVATE
+ lldELF
+ MLIRROCDLToLLVMIRTranslation
+ )
+
+ # Link lldELF also to libmlir.so. Create an alias that starts with LLVM
+ # because LINK_COMPONENTS elements are implicitly prefixed with LLVM.
+ add_library(LLVMAliasTolldELF ALIAS lldELF)
+ set_property(GLOBAL APPEND PROPERTY MLIR_LLVM_LINK_COMPONENTS AliasTolldELF)
+
+endif()
diff --git a/mlir/lib/Dialect/GPU/Transforms/SerializeToHsaco.cpp b/mlir/lib/Dialect/GPU/Transforms/SerializeToHsaco.cpp
new file mode 100644
index 000000000000..1369c1e57549
--- /dev/null
+++ b/mlir/lib/Dialect/GPU/Transforms/SerializeToHsaco.cpp
@@ -0,0 +1,284 @@
+//===- LowerGPUToHSACO.cpp - Convert GPU kernel to HSACO blob -------------===//
+//
+// 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 pass that serializes a gpu module into HSAco blob and
+// adds that blob as a string attribute of the module.
+//
+//===----------------------------------------------------------------------===//
+#include "mlir/Dialect/GPU/Passes.h"
+
+#if MLIR_GPU_TO_HSACO_PASS_ENABLE
+#include "mlir/Pass/Pass.h"
+#include "mlir/Support/FileUtilities.h"
+#include "mlir/Target/LLVMIR/Dialect/ROCDL/ROCDLToLLVMIRTranslation.h"
+#include "mlir/Target/LLVMIR/Export.h"
+
+#include "llvm/MC/MCAsmBackend.h"
+#include "llvm/MC/MCAsmInfo.h"
+#include "llvm/MC/MCCodeEmitter.h"
+#include "llvm/MC/MCContext.h"
+#include "llvm/MC/MCObjectFileInfo.h"
+#include "llvm/MC/MCObjectWriter.h"
+#include "llvm/MC/MCParser/MCTargetAsmParser.h"
+#include "llvm/MC/MCStreamer.h"
+#include "llvm/MC/MCSubtargetInfo.h"
+
+#include "llvm/Support/FileUtilities.h"
+#include "llvm/Support/LineIterator.h"
+#include "llvm/Support/Program.h"
+#include "llvm/Support/TargetRegistry.h"
+#include "llvm/Support/TargetSelect.h"
+#include "llvm/Support/WithColor.h"
+#include "llvm/Target/TargetOptions.h"
+
+#include "lld/Common/Driver.h"
+
+#include "hip/hip_version.h"
+
+#include <mutex>
+
+using namespace mlir;
+
+namespace {
+class SerializeToHsacoPass
+ : public PassWrapper<SerializeToHsacoPass, gpu::SerializeToBlobPass> {
+public:
+ SerializeToHsacoPass();
+
+private:
+ void getDependentDialects(DialectRegistry ®istry) const override;
+
+ // Serializes ROCDL to HSACO.
+ std::unique_ptr<std::vector<char>>
+ serializeISA(const std::string &isa) override;
+
+ std::unique_ptr<SmallVectorImpl<char>> assembleIsa(const std::string &isa);
+ std::unique_ptr<std::vector<char>>
+ createHsaco(const SmallVectorImpl<char> &isaBinary);
+};
+} // namespace
+
+static std::string getDefaultChip() {
+ const char kDefaultChip[] = "gfx900";
+
+ // Locate rocm_agent_enumerator.
+ const char kRocmAgentEnumerator[] = "rocm_agent_enumerator";
+ llvm::ErrorOr<std::string> rocmAgentEnumerator = llvm::sys::findProgramByName(
+ kRocmAgentEnumerator, {__ROCM_PATH__ "/bin"});
+ if (!rocmAgentEnumerator) {
+ llvm::WithColor::warning(llvm::errs())
+ << kRocmAgentEnumerator << "couldn't be located under " << __ROCM_PATH__
+ << "/bin\n";
+ return kDefaultChip;
+ }
+
+ // Prepare temp file to hold the outputs.
+ int tempFd = -1;
+ SmallString<128> tempFilename;
+ if (llvm::sys::fs::createTemporaryFile("rocm_agent", "txt", tempFd,
+ tempFilename)) {
+ llvm::WithColor::warning(llvm::errs())
+ << "temporary file for " << kRocmAgentEnumerator << " creation error\n";
+ return kDefaultChip;
+ }
+ llvm::FileRemover cleanup(tempFilename);
+
+ // Invoke rocm_agent_enumerator.
+ std::string errorMessage;
+ SmallVector<StringRef, 2> args{"-t", "GPU"};
+ Optional<StringRef> redirects[3] = {{""}, tempFilename.str(), {""}};
+ int result =
+ llvm::sys::ExecuteAndWait(rocmAgentEnumerator.get(), args, llvm::None,
+ redirects, 0, 0, &errorMessage);
+ if (result) {
+ llvm::WithColor::warning(llvm::errs())
+ << kRocmAgentEnumerator << " invocation error: " << errorMessage
+ << "\n";
+ return kDefaultChip;
+ }
+
+ // Load and parse the result.
+ auto gfxIsaList = openInputFile(tempFilename);
+ if (!gfxIsaList) {
+ llvm::WithColor::error(llvm::errs())
+ << "read ROCm agent list temp file error\n";
+ return kDefaultChip;
+ }
+ for (llvm::line_iterator lines(*gfxIsaList); !lines.is_at_end(); ++lines) {
+ // Skip the line with content "gfx000".
+ if (*lines == "gfx000")
+ continue;
+ // Use the first ISA version found.
+ return lines->str();
+ }
+
+ return kDefaultChip;
+}
+
+// Sets the 'option' to 'value' unless it already has a value.
+static void maybeSetOption(Pass::Option<std::string> &option,
+ function_ref<std::string()> getValue) {
+ if (!option.hasValue())
+ option = getValue();
+}
+
+SerializeToHsacoPass::SerializeToHsacoPass() {
+ maybeSetOption(this->triple, [] { return "amdgcn-amd-amdhsa"; });
+ maybeSetOption(this->chip, [] {
+ static auto chip = getDefaultChip();
+ return chip;
+ });
+}
+
+void SerializeToHsacoPass::getDependentDialects(
+ DialectRegistry ®istry) const {
+ registerROCDLDialectTranslation(registry);
+ gpu::SerializeToBlobPass::getDependentDialects(registry);
+}
+
+std::unique_ptr<SmallVectorImpl<char>>
+SerializeToHsacoPass::assembleIsa(const std::string &isa) {
+ auto loc = getOperation().getLoc();
+
+ SmallVector<char, 0> result;
+ llvm::raw_svector_ostream os(result);
+
+ llvm::Triple triple(llvm::Triple::normalize(this->triple));
+ std::string error;
+ const llvm::Target *target =
+ llvm::TargetRegistry::lookupTarget(triple.normalize(), error);
+ if (!target) {
+ emitError(loc, Twine("failed to lookup target: ") + error);
+ return {};
+ }
+
+ llvm::SourceMgr srcMgr;
+ srcMgr.AddNewSourceBuffer(llvm::MemoryBuffer::getMemBuffer(isa),
+ llvm::SMLoc());
+
+ const llvm::MCTargetOptions mcOptions;
+ std::unique_ptr<llvm::MCRegisterInfo> mri(
+ target->createMCRegInfo(this->triple));
+ std::unique_ptr<llvm::MCAsmInfo> mai(
+ target->createMCAsmInfo(*mri, this->triple, mcOptions));
+ mai->setRelaxELFRelocations(true);
+
+ llvm::MCObjectFileInfo mofi;
+ llvm::MCContext ctx(mai.get(), mri.get(), &mofi, &srcMgr, &mcOptions);
+ mofi.InitMCObjectFileInfo(triple, false, ctx, false);
+
+ SmallString<128> cwd;
+ if (!llvm::sys::fs::current_path(cwd))
+ ctx.setCompilationDir(cwd);
+
+ std::unique_ptr<llvm::MCStreamer> mcStreamer;
+ std::unique_ptr<llvm::MCInstrInfo> mcii(target->createMCInstrInfo());
+ std::unique_ptr<llvm::MCSubtargetInfo> sti(
+ target->createMCSubtargetInfo(this->triple, this->chip, this->features));
+
+ llvm::MCCodeEmitter *ce = target->createMCCodeEmitter(*mcii, *mri, ctx);
+ llvm::MCAsmBackend *mab = target->createMCAsmBackend(*sti, *mri, mcOptions);
+ mcStreamer.reset(target->createMCObjectStreamer(
+ triple, ctx, std::unique_ptr<llvm::MCAsmBackend>(mab),
+ mab->createObjectWriter(os), std::unique_ptr<llvm::MCCodeEmitter>(ce),
+ *sti, mcOptions.MCRelaxAll, mcOptions.MCIncrementalLinkerCompatible,
+ /*DWARFMustBeAtTheEnd*/ false));
+ mcStreamer->setUseAssemblerInfoForParsing(true);
+
+ std::unique_ptr<llvm::MCAsmParser> parser(
+ createMCAsmParser(srcMgr, ctx, *mcStreamer, *mai));
+ std::unique_ptr<llvm::MCTargetAsmParser> tap(
+ target->createMCAsmParser(*sti, *parser, *mcii, mcOptions));
+
+ if (!tap) {
+ emitError(loc, "assembler initialization error");
+ return {};
+ }
+
+ parser->setTargetParser(*tap);
+ parser->Run(false);
+
+ return std::make_unique<SmallVector<char, 0>>(std::move(result));
+}
+
+std::unique_ptr<std::vector<char>>
+SerializeToHsacoPass::createHsaco(const SmallVectorImpl<char> &isaBinary) {
+ auto loc = getOperation().getLoc();
+
+ // Save the ISA binary to a temp file.
+ int tempIsaBinaryFd = -1;
+ SmallString<128> tempIsaBinaryFilename;
+ if (llvm::sys::fs::createTemporaryFile("kernel", "o", tempIsaBinaryFd,
+ tempIsaBinaryFilename)) {
+ emitError(loc, "temporary file for ISA binary creation error");
+ return {};
+ }
+ llvm::FileRemover cleanupIsaBinary(tempIsaBinaryFilename);
+ llvm::raw_fd_ostream tempIsaBinaryOs(tempIsaBinaryFd, true);
+ tempIsaBinaryOs << StringRef(isaBinary.data(), isaBinary.size());
+ tempIsaBinaryOs.close();
+
+ // Create a temp file for HSA code object.
+ int tempHsacoFD = -1;
+ SmallString<128> tempHsacoFilename;
+ if (llvm::sys::fs::createTemporaryFile("kernel", "hsaco", tempHsacoFD,
+ tempHsacoFilename)) {
+ emitError(loc, "temporary file for HSA code object creation error");
+ return {};
+ }
+ llvm::FileRemover cleanupHsaco(tempHsacoFilename);
+
+ {
+ static std::mutex mutex;
+ const std::lock_guard<std::mutex> lock(mutex);
+ // Invoke lld. Expect a true return value from lld.
+ if (!lld::elf::link({"ld.lld", "-shared", tempIsaBinaryFilename.c_str(),
+ "-o", tempHsacoFilename.c_str()},
+ /*canEarlyExit=*/false, llvm::outs(), llvm::errs())) {
+ emitError(loc, "lld invocation error");
+ return {};
+ }
+ }
+
+ // Load the HSA code object.
+ auto hsacoFile = openInputFile(tempHsacoFilename);
+ if (!hsacoFile) {
+ emitError(loc, "read HSA code object from temp file error");
+ return {};
+ }
+
+ StringRef buffer = hsacoFile->getBuffer();
+ return std::make_unique<std::vector<char>>(buffer.begin(), buffer.end());
+}
+
+std::unique_ptr<std::vector<char>>
+SerializeToHsacoPass::serializeISA(const std::string &isa) {
+ auto isaBinary = assembleIsa(isa);
+ if (!isaBinary)
+ return {};
+ return createHsaco(*isaBinary);
+}
+
+// Register pass to serialize GPU kernel functions to a HSACO binary annotation.
+void mlir::registerGpuSerializeToHsacoPass() {
+ PassRegistration<SerializeToHsacoPass> registerSerializeToHSACO(
+ "gpu-to-hsaco", "Lower GPU kernel function to HSACO binary annotations",
+ [] {
+ // Initialize LLVM AMDGPU backend.
+ LLVMInitializeAMDGPUAsmParser();
+ LLVMInitializeAMDGPUAsmPrinter();
+ LLVMInitializeAMDGPUTarget();
+ LLVMInitializeAMDGPUTargetInfo();
+ LLVMInitializeAMDGPUTargetMC();
+
+ return std::make_unique<SerializeToHsacoPass>();
+ });
+}
+#else // MLIR_GPU_TO_HSACO_PASS_ENABLE
+void mlir::registerGpuSerializeToHsacoPass() {}
+#endif // MLIR_GPU_TO_HSACO_PASS_ENABLE
diff --git a/mlir/lib/ExecutionEngine/CMakeLists.txt b/mlir/lib/ExecutionEngine/CMakeLists.txt
index b9176cf1e89b..978bf1adedd5 100644
--- a/mlir/lib/ExecutionEngine/CMakeLists.txt
+++ b/mlir/lib/ExecutionEngine/CMakeLists.txt
@@ -7,6 +7,7 @@ set(LLVM_OPTIONAL_SOURCES
CudaRuntimeWrappers.cpp
SparseUtils.cpp
ExecutionEngine.cpp
+ RocmRuntimeWrappers.cpp
RunnerUtils.cpp
OptUtils.cpp
JitRunner.cpp
@@ -136,3 +137,51 @@ if(MLIR_CUDA_RUNNER_ENABLED)
${CUDA_RUNTIME_LIBRARY}
)
endif()
+
+if(MLIR_ROCM_RUNNER_ENABLED)
+ # Configure ROCm support.
+ if (NOT DEFINED ROCM_PATH)
+ if (NOT DEFINED ENV{ROCM_PATH})
+ set(ROCM_PATH "/opt/rocm" CACHE PATH "Path to which ROCm has been installed")
+ else()
+ set(ROCM_PATH $ENV{ROCM_PATH} CACHE PATH "Path to which ROCm has been installed")
+ endif()
+ set(HIP_PATH "${ROCM_PATH}/hip" CACHE PATH "Path to which HIP has been installed")
+ endif()
+ set(CMAKE_MODULE_PATH "${HIP_PATH}/cmake" ${CMAKE_MODULE_PATH})
+ find_package(HIP)
+ if (NOT HIP_FOUND)
+ message(SEND_ERROR "Building mlir with ROCm support requires a working ROCm and HIP install")
+ else()
+ message(STATUS "ROCm HIP version: ${HIP_VERSION}")
+ endif()
+
+ # Locate HIP runtime library.
+ find_library(ROCM_RUNTIME_LIBRARY amdhip64
+ PATHS "${HIP_PATH}/lib")
+ if (NOT ROCM_RUNTIME_LIBRARY)
+ message(SEND_ERROR "Could not locate ROCm HIP runtime library")
+ else()
+ message(STATUS "ROCm HIP runtime lib: ${ROCM_RUNTIME_LIBRARY}")
+ endif()
+
+ add_mlir_library(mlir_rocm_runtime
+ SHARED
+ RocmRuntimeWrappers.cpp
+
+ EXCLUDE_FROM_LIBMLIR
+ )
+ target_compile_definitions(mlir_rocm_runtime
+ PRIVATE
+ __HIP_PLATFORM_HCC__
+ )
+ target_include_directories(mlir_rocm_runtime
+ PRIVATE
+ ${HIP_PATH}/include
+ ${ROCM_PATH}/include
+ )
+ target_link_libraries(mlir_rocm_runtime
+ PRIVATE
+ ${ROCM_RUNTIME_LIBRARY}
+ )
+endif()
diff --git a/mlir/tools/mlir-rocm-runner/rocm-runtime-wrappers.cpp b/mlir/lib/ExecutionEngine/RocmRuntimeWrappers.cpp
similarity index 91%
rename from mlir/tools/mlir-rocm-runner/rocm-runtime-wrappers.cpp
rename to mlir/lib/ExecutionEngine/RocmRuntimeWrappers.cpp
index 361ba8f8529d..399a37331060 100644
--- a/mlir/tools/mlir-rocm-runner/rocm-runtime-wrappers.cpp
+++ b/mlir/lib/ExecutionEngine/RocmRuntimeWrappers.cpp
@@ -1,4 +1,4 @@
-//===- rocm-runtime-wrappers.cpp - MLIR ROCM runner wrapper library -------===//
+//===- RocmRuntimeWrappers.cpp - MLIR ROCM runtime wrapper library --------===//
//
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
// See https://llvm.org/LICENSE.txt for license information.
@@ -30,29 +30,25 @@
fprintf(stderr, "'%s' failed with '%s'\n", #expr, name); \
}(expr)
-// Static reference to HIP primary context for device ordinal 0.
-static hipCtx_t Context = [] {
- HIP_REPORT_IF_ERROR(hipInit(/*flags=*/0));
- hipDevice_t device;
- HIP_REPORT_IF_ERROR(hipDeviceGet(&device, /*ordinal=*/0));
- hipCtx_t context;
- HIP_REPORT_IF_ERROR(hipDevicePrimaryCtxRetain(&context, device));
- return context;
-}();
-
// Sets the `Context` for the duration of the instance and restores the previous
// context on destruction.
class ScopedContext {
public:
ScopedContext() {
- HIP_REPORT_IF_ERROR(hipCtxGetCurrent(&previous));
- HIP_REPORT_IF_ERROR(hipCtxSetCurrent(Context));
+ // Static reference to HIP primary context for device ordinal 0.
+ static hipCtx_t context = [] {
+ HIP_REPORT_IF_ERROR(hipInit(/*flags=*/0));
+ hipDevice_t device;
+ HIP_REPORT_IF_ERROR(hipDeviceGet(&device, /*ordinal=*/0));
+ hipCtx_t ctx;
+ HIP_REPORT_IF_ERROR(hipDevicePrimaryCtxRetain(&ctx, device));
+ return ctx;
+ }();
+
+ HIP_REPORT_IF_ERROR(hipCtxPushCurrent(context));
}
- ~ScopedContext() { HIP_REPORT_IF_ERROR(hipCtxSetCurrent(previous)); }
-
-private:
- hipCtx_t previous;
+ ~ScopedContext() { HIP_REPORT_IF_ERROR(hipCtxPopCurrent(nullptr)); }
};
extern "C" hipModule_t mgpuModuleLoad(void *data) {
diff --git a/mlir/test/CMakeLists.txt b/mlir/test/CMakeLists.txt
index 69d123d02047..775a462db53d 100644
--- a/mlir/test/CMakeLists.txt
+++ b/mlir/test/CMakeLists.txt
@@ -21,8 +21,7 @@ set(MLIR_DIALECT_LINALG_INTEGRATION_TEST_LIB_DIR ${CMAKE_LIBRARY_OUTPUT_DIRECTOR
set(MLIR_RUNNER_UTILS_DIR ${CMAKE_LIBRARY_OUTPUT_DIRECTORY})
# Passed to lit.site.cfg.py.in to set up the path where to find the libraries
-# for the mlir rocm / spirv / vulkan runner tests.
-set(MLIR_ROCM_WRAPPER_LIBRARY_DIR ${CMAKE_LIBRARY_OUTPUT_DIRECTORY})
+# for the mlir spirv / vulkan runner tests.
set(MLIR_SPIRV_WRAPPER_LIBRARY_DIR ${CMAKE_LIBRARY_OUTPUT_DIRECTORY})
set(MLIR_VULKAN_WRAPPER_LIBRARY_DIR ${CMAKE_LIBRARY_OUTPUT_DIRECTORY})
@@ -75,6 +74,10 @@ if(MLIR_CUDA_RUNNER_ENABLED)
list(APPEND MLIR_TEST_DEPENDS mlir_cuda_runtime)
endif()
+if(MLIR_ROCM_RUNNER_ENABLED)
+ list(APPEND MLIR_TEST_DEPENDS mlir_rocm_runtime)
+endif()
+
list(APPEND MLIR_TEST_DEPENDS MLIRUnitTests)
if(LLVM_BUILD_EXAMPLES)
@@ -89,12 +92,6 @@ if(LLVM_BUILD_EXAMPLES)
)
endif()
-if(MLIR_ROCM_RUNNER_ENABLED)
- list(APPEND MLIR_TEST_DEPENDS
- mlir-rocm-runner
- )
-endif()
-
if(MLIR_SPIRV_CPU_RUNNER_ENABLED)
add_subdirectory(mlir-spirv-cpu-runner)
list(APPEND MLIR_TEST_DEPENDS
diff --git a/mlir/test/Conversion/GPUToROCm/lower-rocdl-kernel-to-hsaco.mlir b/mlir/test/Conversion/GPUToROCm/lower-rocdl-kernel-to-hsaco.mlir
index 3d7deb906e77..fb19ac6491b3 100644
--- a/mlir/test/Conversion/GPUToROCm/lower-rocdl-kernel-to-hsaco.mlir
+++ b/mlir/test/Conversion/GPUToROCm/lower-rocdl-kernel-to-hsaco.mlir
@@ -1,6 +1,6 @@
-// RUN: mlir-opt %s --test-kernel-to-hsaco -split-input-file | FileCheck %s
+// RUN: mlir-opt %s --test-gpu-to-hsaco | FileCheck %s
-// CHECK: attributes {rocdl.hsaco = "HSACO"}
+// CHECK: gpu.module @foo attributes {gpu.binary = "HSACO"}
gpu.module @foo {
llvm.func @kernel(%arg0 : f32, %arg1 : !llvm.ptr<f32>)
// CHECK: attributes {gpu.kernel}
@@ -9,8 +9,7 @@ gpu.module @foo {
}
}
-// -----
-
+// CHECK: gpu.module @bar attributes {gpu.binary = "HSACO"}
gpu.module @bar {
// CHECK: func @kernel_a
llvm.func @kernel_a()
diff --git a/mlir/test/Integration/GPU/CUDA/lit.local.cfg b/mlir/test/Integration/GPU/CUDA/lit.local.cfg
index b063ddda7e1d..0bdebfedeee3 100644
--- a/mlir/test/Integration/GPU/CUDA/lit.local.cfg
+++ b/mlir/test/Integration/GPU/CUDA/lit.local.cfg
@@ -1,2 +1,2 @@
if not config.enable_cuda_runner:
- config.unsupported = True
\ No newline at end of file
+ config.unsupported = True
diff --git a/mlir/test/mlir-rocm-runner/gpu-to-hsaco.mlir b/mlir/test/Integration/GPU/ROCM/gpu-to-hsaco.mlir
similarity index 82%
rename from mlir/test/mlir-rocm-runner/gpu-to-hsaco.mlir
rename to mlir/test/Integration/GPU/ROCM/gpu-to-hsaco.mlir
index 3f2d44fca38f..fdc525bd2659 100644
--- a/mlir/test/mlir-rocm-runner/gpu-to-hsaco.mlir
+++ b/mlir/test/Integration/GPU/ROCM/gpu-to-hsaco.mlir
@@ -1,5 +1,9 @@
-// RUN: mlir-rocm-runner %s \
-// RUN: --shared-libs=%cuda_wrapper_library_dir/libcuda-runtime-wrappers%shlibext \
+// RUN: mlir-opt %s \
+// RUN: -gpu-kernel-outlining \
+// RUN: -pass-pipeline='gpu.module(strip-debuginfo,convert-gpu-to-rocdl,gpu-to-hsaco)' \
+// RUN: -gpu-to-llvm \
+// RUN: | mlir-cpu-runner \
+// RUN: --shared-libs=%linalg_test_lib_dir/libmlir_rocm_runtime%shlibext \
// RUN: --shared-libs=%linalg_test_lib_dir/libmlir_runner_utils%shlibext \
// RUN: --entry-point-result=void \
// RUN: | FileCheck %s
diff --git a/mlir/test/mlir-rocm-runner/lit.local.cfg b/mlir/test/Integration/GPU/ROCM/lit.local.cfg
similarity index 100%
rename from mlir/test/mlir-rocm-runner/lit.local.cfg
rename to mlir/test/Integration/GPU/ROCM/lit.local.cfg
diff --git a/mlir/test/mlir-rocm-runner/two-modules.mlir b/mlir/test/Integration/GPU/ROCM/two-modules.mlir
similarity index 74%
rename from mlir/test/mlir-rocm-runner/two-modules.mlir
rename to mlir/test/Integration/GPU/ROCM/two-modules.mlir
index 7c0faae5d135..3c6c56b0091a 100644
--- a/mlir/test/mlir-rocm-runner/two-modules.mlir
+++ b/mlir/test/Integration/GPU/ROCM/two-modules.mlir
@@ -1,5 +1,9 @@
-// RUN: mlir-rocm-runner %s \
-// RUN: --shared-libs=%cuda_wrapper_library_dir/libcuda-runtime-wrappers%shlibext \
+// RUN: mlir-opt %s \
+// RUN: -gpu-kernel-outlining \
+// RUN: -pass-pipeline='gpu.module(strip-debuginfo,convert-gpu-to-rocdl,gpu-to-hsaco)' \
+// RUN: -gpu-to-llvm \
+// RUN: | mlir-cpu-runner \
+// RUN: --shared-libs=%linalg_test_lib_dir/libmlir_rocm_runtime%shlibext \
// RUN: --shared-libs=%linalg_test_lib_dir/libmlir_runner_utils%shlibext \
// RUN: --entry-point-result=void \
// RUN: | FileCheck %s
@@ -30,5 +34,5 @@ func @main() {
return
}
-func @mgpuMemGetDeviceMemRef1dInt32(%ptr : memref<?xi32>) -> (memref<?xi32>)
-func @print_memref_i32(%ptr : memref<*xi32>)
+func private @mgpuMemGetDeviceMemRef1dInt32(%ptr : memref<?xi32>) -> (memref<?xi32>)
+func private @print_memref_i32(%ptr : memref<*xi32>)
diff --git a/mlir/test/mlir-rocm-runner/vecadd.mlir b/mlir/test/Integration/GPU/ROCM/vecadd.mlir
similarity index 87%
rename from mlir/test/mlir-rocm-runner/vecadd.mlir
rename to mlir/test/Integration/GPU/ROCM/vecadd.mlir
index d4dc862c60b6..917be3c93d08 100644
--- a/mlir/test/mlir-rocm-runner/vecadd.mlir
+++ b/mlir/test/Integration/GPU/ROCM/vecadd.mlir
@@ -1,5 +1,10 @@
-// RUN: mlir-rocm-runner %s \
-// RUN: --shared-libs=%cuda_wrapper_library_dir/libcuda-runtime-wrappers%shlibext \
+// RUN: mlir-opt %s \
+// RUN: -convert-scf-to-std \
+// RUN: -gpu-kernel-outlining \
+// RUN: -pass-pipeline='gpu.module(strip-debuginfo,convert-gpu-to-rocdl,gpu-to-hsaco)' \
+// RUN: -gpu-to-llvm \
+// RUN: | mlir-cpu-runner \
+// RUN: --shared-libs=%linalg_test_lib_dir/libmlir_rocm_runtime%shlibext \
// RUN: --shared-libs=%linalg_test_lib_dir/libmlir_runner_utils%shlibext \
// RUN: --entry-point-result=void \
// RUN: | FileCheck %s
diff --git a/mlir/test/mlir-rocm-runner/vector-transferops.mlir b/mlir/test/Integration/GPU/ROCM/vector-transferops.mlir
similarity index 90%
rename from mlir/test/mlir-rocm-runner/vector-transferops.mlir
rename to mlir/test/Integration/GPU/ROCM/vector-transferops.mlir
index eda541a2d814..c2807b64c064 100644
--- a/mlir/test/mlir-rocm-runner/vector-transferops.mlir
+++ b/mlir/test/Integration/GPU/ROCM/vector-transferops.mlir
@@ -1,5 +1,10 @@
-// RUN: mlir-rocm-runner %s \
-// RUN: --shared-libs=%cuda_wrapper_library_dir/libcuda-runtime-wrappers%shlibext \
+// RUN: mlir-opt %s \
+// RUN: -convert-scf-to-std \
+// RUN: -gpu-kernel-outlining \
+// RUN: -pass-pipeline='gpu.module(strip-debuginfo,convert-gpu-to-rocdl,gpu-to-hsaco)' \
+// RUN: -gpu-to-llvm \
+// RUN: | mlir-cpu-runner \
+// RUN: --shared-libs=%linalg_test_lib_dir/libmlir_rocm_runtime%shlibext \
// RUN: --shared-libs=%linalg_test_lib_dir/libmlir_runner_utils%shlibext \
// RUN: --entry-point-result=void \
// RUN: | FileCheck %s
diff --git a/mlir/test/lib/Transforms/TestConvertGPUKernelToHsaco.cpp b/mlir/test/lib/Transforms/TestConvertGPUKernelToHsaco.cpp
index 58e890b907a4..5a3cb33526f6 100644
--- a/mlir/test/lib/Transforms/TestConvertGPUKernelToHsaco.cpp
+++ b/mlir/test/lib/Transforms/TestConvertGPUKernelToHsaco.cpp
@@ -6,11 +6,9 @@
//
//===----------------------------------------------------------------------===//
-#include "mlir/Conversion/GPUCommon/GPUCommonPass.h"
-#include "mlir/Dialect/LLVMIR/ROCDLDialect.h"
+#include "mlir/Dialect/GPU/Passes.h"
+
#include "mlir/Pass/Pass.h"
-#include "mlir/Pass/PassManager.h"
-#include "mlir/Target/LLVMIR/Dialect/LLVMIR/LLVMToLLVMIRTranslation.h"
#include "mlir/Target/LLVMIR/Dialect/ROCDL/ROCDLToLLVMIRTranslation.h"
#include "mlir/Target/LLVMIR/Export.h"
#include "llvm/Support/TargetSelect.h"
@@ -18,38 +16,54 @@
using namespace mlir;
#if MLIR_ROCM_CONVERSIONS_ENABLED
-static OwnedBlob compileIsaToHsacoForTesting(const std::string &, Location,
- StringRef) {
- const char data[] = "HSACO";
- return std::make_unique<std::vector<char>>(data, data + sizeof(data) - 1);
+namespace {
+class TestSerializeToHsacoPass
+ : public PassWrapper<TestSerializeToHsacoPass, gpu::SerializeToBlobPass> {
+public:
+ TestSerializeToHsacoPass();
+
+private:
+ void getDependentDialects(DialectRegistry ®istry) const override;
+
+ // Serializes ROCDL IR to HSACO.
+ std::unique_ptr<std::vector<char>>
+ serializeISA(const std::string &isa) override;
+};
+} // namespace
+
+TestSerializeToHsacoPass::TestSerializeToHsacoPass() {
+ this->triple = "amdgcn-amd-amdhsa";
+ this->chip = "gfx900";
+}
+
+void TestSerializeToHsacoPass::getDependentDialects(
+ DialectRegistry ®istry) const {
+ registerROCDLDialectTranslation(registry);
+ gpu::SerializeToBlobPass::getDependentDialects(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);
+std::unique_ptr<std::vector<char>>
+TestSerializeToHsacoPass::serializeISA(const std::string &) {
+ std::string data = "HSACO";
+ return std::make_unique<std::vector<char>>(data.begin(), data.end());
}
namespace mlir {
namespace test {
-void registerTestConvertGPUKernelToHsacoPass() {
- PassPipelineRegistration<>(
- "test-kernel-to-hsaco",
- "Convert all kernel functions to ROCm hsaco blobs",
- [](OpPassManager &pm) {
+// Register test pass to serialize GPU module to a HSAco binary annotation.
+void registerTestGpuSerializeToHsacoPass() {
+ PassRegistration<TestSerializeToHsacoPass> registerSerializeToHsaco(
+ "test-gpu-to-hsaco",
+ "Lower GPU kernel function to HSAco binary annotations", [] {
// Initialize LLVM AMDGPU backend.
LLVMInitializeAMDGPUTarget();
LLVMInitializeAMDGPUTargetInfo();
LLVMInitializeAMDGPUTargetMC();
LLVMInitializeAMDGPUAsmPrinter();
- pm.addPass(createConvertGPUKernelToBlobPass(
- translateModuleToROCDL, compileIsaToHsacoForTesting,
- "amdgcn-amd-amdhsa", "gfx900", "-code-object-v3", "rocdl.hsaco"));
+ return std::make_unique<TestSerializeToHsacoPass>();
});
}
} // namespace test
} // namespace mlir
-#endif
+#endif // MLIR_ROCM_CONVERSIONS_ENABLED
diff --git a/mlir/test/lit.cfg.py b/mlir/test/lit.cfg.py
index 4ba36202578d..199d7222e1cf 100644
--- a/mlir/test/lit.cfg.py
+++ b/mlir/test/lit.cfg.py
@@ -77,7 +77,6 @@
ToolSubst('toy-ch5', unresolved='ignore'),
ToolSubst('%linalg_test_lib_dir', config.linalg_test_lib_dir, unresolved='ignore'),
ToolSubst('%mlir_runner_utils_dir', config.mlir_runner_utils_dir, unresolved='ignore'),
- ToolSubst('%rocm_wrapper_library_dir', config.rocm_wrapper_library_dir, unresolved='ignore'),
ToolSubst('%spirv_wrapper_library_dir', config.spirv_wrapper_library_dir, unresolved='ignore'),
ToolSubst('%vulkan_wrapper_library_dir', config.vulkan_wrapper_library_dir, unresolved='ignore'),
ToolSubst('%mlir_integration_test_dir', config.mlir_integration_test_dir, unresolved='ignore'),
diff --git a/mlir/test/lit.site.cfg.py.in b/mlir/test/lit.site.cfg.py.in
index 0015c1369d7a..dbc8460df576 100644
--- a/mlir/test/lit.site.cfg.py.in
+++ b/mlir/test/lit.site.cfg.py.in
@@ -39,7 +39,6 @@ config.build_examples = @LLVM_BUILD_EXAMPLES@
config.run_cuda_tests = @MLIR_CUDA_CONVERSIONS_ENABLED@
config.enable_cuda_runner = @MLIR_CUDA_RUNNER_ENABLED@
config.run_rocm_tests = @MLIR_ROCM_CONVERSIONS_ENABLED@
-config.rocm_wrapper_library_dir = "@MLIR_ROCM_WRAPPER_LIBRARY_DIR@"
config.enable_rocm_runner = @MLIR_ROCM_RUNNER_ENABLED@
config.spirv_wrapper_library_dir = "@MLIR_SPIRV_WRAPPER_LIBRARY_DIR@"
config.enable_spirv_cpu_runner = @MLIR_SPIRV_CPU_RUNNER_ENABLED@
diff --git a/mlir/tools/CMakeLists.txt b/mlir/tools/CMakeLists.txt
index 37793ce65ab1..ac9ca8167320 100644
--- a/mlir/tools/CMakeLists.txt
+++ b/mlir/tools/CMakeLists.txt
@@ -1,7 +1,6 @@
add_subdirectory(mlir-cpu-runner)
add_subdirectory(mlir-opt)
add_subdirectory(mlir-reduce)
-add_subdirectory(mlir-rocm-runner)
add_subdirectory(mlir-shlib)
add_subdirectory(mlir-spirv-cpu-runner)
add_subdirectory(mlir-translate)
diff --git a/mlir/tools/mlir-opt/mlir-opt.cpp b/mlir/tools/mlir-opt/mlir-opt.cpp
index 241cee572cb1..428b3d506317 100644
--- a/mlir/tools/mlir-opt/mlir-opt.cpp
+++ b/mlir/tools/mlir-opt/mlir-opt.cpp
@@ -65,7 +65,7 @@ void registerTestCallGraphPass();
void registerTestConstantFold();
void registerTestConvVectorization();
void registerTestGpuSerializeToCubinPass();
-void registerTestConvertGPUKernelToHsacoPass();
+void registerTestGpuSerializeToHsacoPass();
void registerTestDataLayoutQuery();
void registerTestDecomposeCallGraphTypes();
void registerTestDialect(DialectRegistry &);
@@ -140,7 +140,7 @@ void registerTestPasses() {
test::registerTestGpuSerializeToCubinPass();
#endif
#if MLIR_ROCM_CONVERSIONS_ENABLED
- test::registerTestConvertGPUKernelToHsacoPass();
+ test::registerTestGpuSerializeToHsacoPass();
#endif
test::registerTestConvVectorization();
test::registerTestDecomposeCallGraphTypes();
diff --git a/mlir/tools/mlir-rocm-runner/CMakeLists.txt b/mlir/tools/mlir-rocm-runner/CMakeLists.txt
deleted file mode 100644
index d2381413e158..000000000000
--- a/mlir/tools/mlir-rocm-runner/CMakeLists.txt
+++ /dev/null
@@ -1,127 +0,0 @@
-set(LLVM_OPTIONAL_SOURCES
- rocm-runtime-wrappers.cpp
- mlir-rocm-runner.cpp
- )
-
-if(MLIR_ROCM_RUNNER_ENABLED)
- if (NOT ("AMDGPU" IN_LIST LLVM_TARGETS_TO_BUILD))
- message(SEND_ERROR
- "Building the mlir rocm runner requires the AMDGPU backend")
- endif()
-
- # Ensure lld is enabled.
- if (NOT "lld" IN_LIST LLVM_ENABLE_PROJECTS)
- message(SEND_ERROR "lld is not enabled. Please revise LLVM_ENABLE_PROJECTS")
- endif()
-
- # lld header files.
- include_directories(${MLIR_SOURCE_DIR}/../lld/include)
-
- # Configure ROCm support.
- if (NOT DEFINED ROCM_PATH)
- if (NOT DEFINED ENV{ROCM_PATH})
- set(ROCM_PATH "/opt/rocm" CACHE PATH "Path to which ROCm has been installed")
- else()
- set(ROCM_PATH $ENV{ROCM_PATH} CACHE PATH "Path to which ROCm has been installed")
- endif()
- set(HIP_PATH "${ROCM_PATH}/hip" CACHE PATH " Path to which HIP has been installed")
- endif()
- set(CMAKE_MODULE_PATH "${HIP_PATH}/cmake" ${CMAKE_MODULE_PATH})
- find_package(HIP)
- if (NOT HIP_FOUND)
- message(SEND_ERROR "Build the mlir rocm runner requires a working ROCm and HIP install")
- else()
- message(STATUS "ROCm HIP version: ${HIP_VERSION}")
- endif()
-
- # Set compile-time flags for ROCm path.
- add_definitions(-D__ROCM_PATH__="${ROCM_PATH}")
-
- # Locate HIP runtime library.
- find_library(ROCM_RUNTIME_LIBRARY amdhip64
- PATHS "${HIP_PATH}/lib")
- if (NOT ROCM_RUNTIME_LIBRARY)
- message(SEND_ERROR "Could not locate ROCm HIP runtime library")
- else()
- message(STATUS "ROCm HIP runtime lib: ${ROCM_RUNTIME_LIBRARY}")
- endif()
-
- # Set HIP compile-time flags.
- add_definitions(-D__HIP_PLATFORM_HCC__)
-
- add_mlir_library(rocm-runtime-wrappers
- SHARED
- rocm-runtime-wrappers.cpp
-
- EXCLUDE_FROM_LIBMLIR
- )
- target_include_directories(rocm-runtime-wrappers
- PRIVATE
- "${HIP_PATH}/../include"
- "${HIP_PATH}/include"
- )
- target_link_libraries(rocm-runtime-wrappers
- PRIVATE
- ${ROCM_RUNTIME_LIBRARY}
- )
-
- get_property(conversion_libs GLOBAL PROPERTY MLIR_CONVERSION_LIBS)
- set(LIBS
- ${conversion_libs}
- lldCommon
- lldDriver
- lldELF
- MLIRJitRunner
- MLIRAnalysis
- MLIREDSC
- MLIRExecutionEngine
- MLIRGPU
- MLIRIR
- MLIRLLVMIR
- MLIRLLVMToLLVMIRTranslation
- MLIRParser
- MLIRROCDLIR
- MLIRStandard
- MLIRSupport
- MLIRTargetLLVMIRExport
- MLIRROCDLToLLVMIRTranslation
- MLIRTransforms
- MLIRTranslation
- ${ROCM_RUNTIME_LIBRARY}
- )
-
- # Manually expand the target library, since our MLIR libraries
- # aren't plugged into the LLVM dependency tracking. If we don't
- # do this then we can't insert the CodeGen library after ourselves
- llvm_expand_pseudo_components(TARGET_LIBS AllTargetsCodeGens AllTargetsAsmParsers)
- # Prepend LLVM in front of every target, this is how the library
- # are named with CMake
- SET(targets_to_link)
- FOREACH(t ${TARGET_LIBS})
- LIST(APPEND targets_to_link "LLVM${t}")
- ENDFOREACH(t)
-
- add_llvm_tool(mlir-rocm-runner
- mlir-rocm-runner.cpp
-
- DEPENDS
- rocm-runtime-wrappers
-
- LINK_COMPONENTS
-
- Core
- LTO
- MC
- MCParser
- Option
- Support
- )
- llvm_update_compile_flags(mlir-rocm-runner)
- target_include_directories(mlir-rocm-runner
- PRIVATE
- "${HIP_PATH}/../include"
- "${HIP_PATH}/include"
- )
- target_link_libraries(mlir-rocm-runner PRIVATE ${LIBS} ${targets_to_link})
-
-endif()
diff --git a/mlir/tools/mlir-rocm-runner/mlir-rocm-runner.cpp b/mlir/tools/mlir-rocm-runner/mlir-rocm-runner.cpp
deleted file mode 100644
index c2f9abbf73c8..000000000000
--- a/mlir/tools/mlir-rocm-runner/mlir-rocm-runner.cpp
+++ /dev/null
@@ -1,349 +0,0 @@
-//===- mlir-rocm-runner.cpp - MLIR ROCM Execution Driver-------------------===//
-//
-// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
-// See https://llvm.org/LICENSE.txt for license information.
-// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
-//
-//===----------------------------------------------------------------------===//
-//
-// This is a command line utility that executes an MLIR file on the GPU by
-// translating MLIR to ROCDL/LLVM IR before JIT-compiling and executing the
-// latter.
-//
-//===----------------------------------------------------------------------===//
-
-#include "llvm/ADT/STLExtras.h"
-
-#include "mlir/Conversion/GPUCommon/GPUCommonPass.h"
-#include "mlir/Conversion/GPUToROCDL/GPUToROCDLPass.h"
-#include "mlir/Conversion/SCFToStandard/SCFToStandard.h"
-#include "mlir/Conversion/StandardToLLVM/ConvertStandardToLLVM.h"
-#include "mlir/Conversion/StandardToLLVM/ConvertStandardToLLVMPass.h"
-#include "mlir/Dialect/GPU/GPUDialect.h"
-#include "mlir/Dialect/GPU/Passes.h"
-#include "mlir/Dialect/LLVMIR/LLVMDialect.h"
-#include "mlir/Dialect/LLVMIR/ROCDLDialect.h"
-#include "mlir/Dialect/StandardOps/IR/Ops.h"
-#include "mlir/ExecutionEngine/JitRunner.h"
-#include "mlir/ExecutionEngine/OptUtils.h"
-#include "mlir/IR/BuiltinOps.h"
-#include "mlir/Pass/Pass.h"
-#include "mlir/Pass/PassManager.h"
-#include "mlir/Support/FileUtilities.h"
-#include "mlir/Target/LLVMIR/Dialect/LLVMIR/LLVMToLLVMIRTranslation.h"
-#include "mlir/Target/LLVMIR/Dialect/ROCDL/ROCDLToLLVMIRTranslation.h"
-#include "mlir/Target/LLVMIR/Export.h"
-#include "mlir/Transforms/DialectConversion.h"
-#include "mlir/Transforms/Passes.h"
-#include "llvm/Support/ErrorOr.h"
-#include "llvm/Support/FileUtilities.h"
-#include "llvm/Support/InitLLVM.h"
-#include "llvm/Support/LineIterator.h"
-#include "llvm/Support/Program.h"
-#include "llvm/Support/SourceMgr.h"
-#include "llvm/Support/TargetRegistry.h"
-#include "llvm/Support/TargetSelect.h"
-
-// MC headers.
-#include "llvm/MC/MCAsmBackend.h"
-#include "llvm/MC/MCAsmInfo.h"
-#include "llvm/MC/MCCodeEmitter.h"
-#include "llvm/MC/MCContext.h"
-#include "llvm/MC/MCInstPrinter.h"
-#include "llvm/MC/MCInstrInfo.h"
-#include "llvm/MC/MCObjectFileInfo.h"
-#include "llvm/MC/MCObjectWriter.h"
-#include "llvm/MC/MCParser/AsmLexer.h"
-#include "llvm/MC/MCParser/MCTargetAsmParser.h"
-#include "llvm/MC/MCRegisterInfo.h"
-#include "llvm/MC/MCStreamer.h"
-#include "llvm/MC/MCSubtargetInfo.h"
-#include "llvm/MC/MCTargetOptionsCommandFlags.h"
-
-// lld headers.
-#include "lld/Common/Driver.h"
-
-// HIP headers.
-#include "hip/hip_version.h"
-
-#include <mutex>
-
-using namespace mlir;
-using namespace llvm;
-
-using Blob = SmallVector<char, 0>;
-
-static cl::opt<std::string> tripleName("triple", cl::desc("target triple"),
- cl::value_desc("triple string"),
- cl::init("amdgcn-amd-amdhsa"));
-
-static cl::opt<std::string> targetChip("target", cl::desc("target chip"),
- cl::value_desc("AMDGPU ISA version"),
- cl::init(""));
-
-static cl::opt<std::string> features("feature", cl::desc("target features"),
- cl::value_desc("AMDGPU target features"),
- cl::init(""));
-
-static constexpr const char kRunnerProgram[] = "mlir-rocm-runner";
-static constexpr const char kRocmAgentEnumerator[] = "rocm_agent_enumerator";
-static constexpr const char kDefaultTargetChip[] = "gfx900";
-
-static LogicalResult assembleIsa(const std::string isa, StringRef name,
- Blob &result) {
- raw_svector_ostream os(result);
-
- std::string error;
- Triple theTriple(Triple::normalize(tripleName));
- const Target *theTarget =
- TargetRegistry::lookupTarget(theTriple.normalize(), error);
- if (!theTarget) {
- WithColor::error(errs(), name) << error;
- return failure();
- }
-
- SourceMgr srcMgr;
- srcMgr.AddNewSourceBuffer(MemoryBuffer::getMemBuffer(isa), SMLoc());
-
- const MCTargetOptions mcOptions;
- std::unique_ptr<MCRegisterInfo> mri(theTarget->createMCRegInfo(tripleName));
- std::unique_ptr<MCAsmInfo> mai(
- theTarget->createMCAsmInfo(*mri, tripleName, mcOptions));
- mai->setRelaxELFRelocations(true);
-
- MCObjectFileInfo mofi;
- MCContext ctx(mai.get(), mri.get(), &mofi, &srcMgr, &mcOptions);
- mofi.InitMCObjectFileInfo(theTriple, false, ctx, false);
-
- SmallString<128> cwd;
- if (!sys::fs::current_path(cwd))
- ctx.setCompilationDir(cwd);
-
- std::unique_ptr<MCStreamer> mcStreamer;
- std::unique_ptr<MCInstrInfo> mcii(theTarget->createMCInstrInfo());
- std::unique_ptr<MCSubtargetInfo> sti(
- theTarget->createMCSubtargetInfo(tripleName, targetChip, features));
-
- MCCodeEmitter *ce = theTarget->createMCCodeEmitter(*mcii, *mri, ctx);
- MCAsmBackend *mab = theTarget->createMCAsmBackend(*sti, *mri, mcOptions);
- mcStreamer.reset(theTarget->createMCObjectStreamer(
- theTriple, ctx, std::unique_ptr<MCAsmBackend>(mab),
- mab->createObjectWriter(os), std::unique_ptr<MCCodeEmitter>(ce), *sti,
- mcOptions.MCRelaxAll, mcOptions.MCIncrementalLinkerCompatible,
- /*DWARFMustBeAtTheEnd*/ false));
- mcStreamer->setUseAssemblerInfoForParsing(true);
-
- std::unique_ptr<MCAsmParser> parser(
- createMCAsmParser(srcMgr, ctx, *mcStreamer, *mai));
- std::unique_ptr<MCTargetAsmParser> tap(
- theTarget->createMCAsmParser(*sti, *parser, *mcii, mcOptions));
-
- if (!tap) {
- WithColor::error(errs(), name) << "assembler initialization error.\n";
- return failure();
- }
-
- parser->setTargetParser(*tap);
- parser->Run(false);
-
- return success();
-}
-
-static std::mutex mutex;
-static LogicalResult createHsaco(const Blob &isaBlob, StringRef name,
- Blob &hsacoBlob) {
- // Save the ISA binary to a temp file.
- int tempIsaBinaryFd = -1;
- SmallString<128> tempIsaBinaryFilename;
- std::error_code ec = sys::fs::createTemporaryFile(
- "kernel", "o", tempIsaBinaryFd, tempIsaBinaryFilename);
- if (ec) {
- WithColor::error(errs(), name)
- << "temporary file for ISA binary creation error.\n";
- return failure();
- }
- FileRemover cleanupIsaBinary(tempIsaBinaryFilename);
- raw_fd_ostream tempIsaBinaryOs(tempIsaBinaryFd, true);
- tempIsaBinaryOs << isaBlob;
- tempIsaBinaryOs.close();
-
- // Create a temp file for HSA code object.
- int tempHsacoFD = -1;
- SmallString<128> tempHsacoFilename;
- ec = sys::fs::createTemporaryFile("kernel", "hsaco", tempHsacoFD,
- tempHsacoFilename);
- if (ec) {
- WithColor::error(errs(), name)
- << "temporary file for HSA code object creation error.\n";
- return failure();
- }
- FileRemover cleanupHsaco(tempHsacoFilename);
-
- const std::lock_guard<std::mutex> lock(mutex);
- // Invoke lld. Expect a true return value from lld.
- bool ret = lld::elf::link({"ld.lld", "-shared", tempIsaBinaryFilename.c_str(),
- "-o", tempHsacoFilename.c_str()},
- /*canEarlyExit=*/false, llvm::outs(), llvm::errs());
- if (!ret) {
- WithColor::error(errs(), name) << "lld invocation error.\n";
- return failure();
- }
-
- // Load the HSA code object.
- auto hsacoFile = mlir::openInputFile(tempHsacoFilename);
- if (!hsacoFile) {
- WithColor::error(errs(), name)
- << "read HSA code object from temp file error.\n";
- return failure();
- }
- hsacoBlob.assign(hsacoFile->getBuffer().begin(),
- hsacoFile->getBuffer().end());
-
- return success();
-}
-
-static std::unique_ptr<llvm::Module>
-compileModuleToROCDLIR(Operation *m, llvm::LLVMContext &llvmContext,
- StringRef name) {
- auto llvmModule = translateModuleToROCDLIR(m, llvmContext, name);
- // TODO: Link with ROCm-Device-Libs in case needed (ex: the Module
- // depends on math functions).
- return llvmModule;
-}
-
-static OwnedBlob compileISAToHsaco(const std::string isa, Location loc,
- StringRef name) {
- // ISA -> ISA in binary form via MC.
- // Use lld to create HSA code object.
- Blob isaBlob;
- Blob hsacoBlob;
-
- if (succeeded(assembleIsa(isa, name, isaBlob)) &&
- succeeded(createHsaco(isaBlob, name, hsacoBlob)))
- return std::make_unique<std::vector<char>>(hsacoBlob.begin(),
- hsacoBlob.end());
-
- WithColor::error(errs(), name) << "producing HSA code object error.\n";
- return {};
-}
-
-static void configTargetChip() {
- // Set targetChip to default value first.
- targetChip = kDefaultTargetChip;
-
- // Locate rocm_agent_enumerator.
- llvm::ErrorOr<std::string> rocmAgentEnumerator = llvm::sys::findProgramByName(
- kRocmAgentEnumerator, {__ROCM_PATH__ "/bin"});
- std::error_code ec;
- if ((ec = rocmAgentEnumerator.getError())) {
- WithColor::warning(errs(), kRunnerProgram)
- << kRocmAgentEnumerator << " couldn't be located under "
- << __ROCM_PATH__ << ", set target as " << kDefaultTargetChip << "\n";
- return;
- }
-
- // Prepare temp file to hold the outputs.
- int tempFd = -1;
- SmallString<128> tempFilename;
- ec = sys::fs::createTemporaryFile("rocm_agent", "txt", tempFd, tempFilename);
- if (ec) {
- WithColor::warning(errs(), kRunnerProgram)
- << "temporary file for " << kRocmAgentEnumerator
- << " creation error, set target as " << kDefaultTargetChip << "\n";
- return;
- }
- FileRemover cleanup(tempFilename);
-
- // Invoke rocm_agent_enumerator.
- std::string errorMessage;
- SmallVector<StringRef, 2> args{"-t", "GPU"};
- Optional<StringRef> redirects[3] = {{""}, tempFilename.str(), {""}};
- int result =
- llvm::sys::ExecuteAndWait(rocmAgentEnumerator.get(), args, llvm::None,
- redirects, 0, 0, &errorMessage);
- if (result) {
- WithColor::warning(errs(), kRunnerProgram)
- << kRocmAgentEnumerator << " invocation error: " << errorMessage
- << ", set target as " << kDefaultTargetChip << "\n";
- return;
- }
-
- // Load and parse the result.
- auto gfxIsaList = mlir::openInputFile(tempFilename);
- if (!gfxIsaList) {
- WithColor::error(errs(), kRunnerProgram)
- << "read ROCm agent list temp file error, set target as "
- << kDefaultTargetChip << "\n";
- return;
- }
- for (line_iterator lines(*gfxIsaList); !lines.is_at_end(); ++lines) {
- // Skip the line with content "gfx000".
- if (*lines == "gfx000")
- continue;
- // Use the first ISA version found.
- targetChip = lines->str();
- break;
- }
-}
-
-static void configTargetFeatures() {
- if (features.size() > 0)
- features += ",";
- // After ROCm 3.5, adopt HSA code object V3.
- if (HIP_VERSION_MAJOR >= 3 && HIP_VERSION_MINOR >= 5)
- features += "+code-object-v3";
- else
- features += "-code-object-v3";
-}
-
-static LogicalResult runMLIRPasses(ModuleOp m) {
- PassManager pm(m.getContext());
- applyPassManagerCLOptions(pm);
-
- // Configure target chip ISA version if it has not been specified.
- if (!targetChip.size())
- configTargetChip();
-
- // Configure target features per ROCm / HIP version.
- configTargetFeatures();
-
- const char gpuBinaryAnnotation[] = "rocdl.hsaco";
- pm.addPass(createLowerToCFGPass());
- pm.addPass(createGpuKernelOutliningPass());
- auto &kernelPm = pm.nest<gpu::GPUModuleOp>();
- kernelPm.addPass(createStripDebugInfoPass());
- kernelPm.addPass(createLowerGpuOpsToROCDLOpsPass());
- kernelPm.addPass(createConvertGPUKernelToBlobPass(
- compileModuleToROCDLIR, compileISAToHsaco, tripleName, targetChip,
- features, gpuBinaryAnnotation));
- pm.addPass(createGpuToLLVMConversionPass(gpuBinaryAnnotation));
-
- return pm.run(m);
-}
-
-int main(int argc, char **argv) {
- registerPassManagerCLOptions();
- llvm::InitLLVM y(argc, argv);
- llvm::InitializeAllTargetInfos();
- llvm::InitializeAllTargetMCs();
- llvm::InitializeAllAsmParsers();
-
- // Initialize LLVM AMDGPU backend.
- LLVMInitializeAMDGPUTarget();
- LLVMInitializeAMDGPUTargetInfo();
- LLVMInitializeAMDGPUTargetMC();
- LLVMInitializeAMDGPUAsmPrinter();
-
- mlir::initializeLLVMPasses();
-
- mlir::JitRunnerConfig jitRunnerConfig;
- jitRunnerConfig.mlirTransformer = runMLIRPasses;
-
- mlir::DialectRegistry registry;
- registry.insert<mlir::LLVM::LLVMDialect, mlir::gpu::GPUDialect,
- mlir::ROCDL::ROCDLDialect, mlir::StandardOpsDialect>();
- mlir::registerLLVMDialectTranslation(registry);
- mlir::registerROCDLDialectTranslation(registry);
-
- return mlir::JitRunnerMain(argc, argv, registry, jitRunnerConfig);
-}
More information about the Mlir-commits
mailing list