[Mlir-commits] [mlir] [mlir] Remove the mlir-spirv-cpu-runner (move to mlir-cpu-runner) (PR #114563)
llvmlistbot at llvm.org
llvmlistbot at llvm.org
Fri Nov 1 09:23:34 PDT 2024
llvmbot wrote:
<!--LLVM PR SUMMARY COMMENT-->
@llvm/pr-subscribers-mlir-core
Author: Andrea Faulds (andfau-amd)
<details>
<summary>Changes</summary>
This commit builds on and completes the work done in 9f6c632ecda08bfff76b798c46d5d7cfde57b5e9 to eliminate the need for a separate mlir-spirv-cpu-runner binary. Since the MLIR processing is already done outside this runner, the only real difference between it and the mlir-cpu-runner is the final linking step between the nested LLVM IR modules. By moving this step into mlir-cpu-runner behind a new command-line flag (`--link-nested-modules`), this commit is able to completely remvoe the runner component of the mlir-spirv-cpu-runner. The runtime libraries and associated build infrastructure are left undisturbed.
---
Full diff: https://github.com/llvm/llvm-project/pull/114563.diff
8 Files Affected:
- (modified) mlir/docs/SPIRVToLLVMDialectConversion.md (+6-4)
- (modified) mlir/test/CMakeLists.txt (-1)
- (modified) mlir/test/mlir-spirv-cpu-runner/double.mlir (+1-1)
- (modified) mlir/test/mlir-spirv-cpu-runner/simple_add.mlir (+1-1)
- (modified) mlir/tools/CMakeLists.txt (-1)
- (modified) mlir/tools/mlir-cpu-runner/mlir-cpu-runner.cpp (+58-1)
- (removed) mlir/tools/mlir-spirv-cpu-runner/CMakeLists.txt (-36)
- (removed) mlir/tools/mlir-spirv-cpu-runner/mlir-spirv-cpu-runner.cpp (-90)
``````````diff
diff --git a/mlir/docs/SPIRVToLLVMDialectConversion.md b/mlir/docs/SPIRVToLLVMDialectConversion.md
index 0aae02cff26be1..d77a86de8a6846 100644
--- a/mlir/docs/SPIRVToLLVMDialectConversion.md
+++ b/mlir/docs/SPIRVToLLVMDialectConversion.md
@@ -817,12 +817,14 @@ to LLVM ops. At the moment, SPIR-V module attributes are ignored.
## `mlir-spirv-cpu-runner`
-`mlir-spirv-cpu-runner` allows to execute `gpu` dialect kernel on the CPU via
-SPIR-V to LLVM dialect conversion. Currently, only single-threaded kernel is
+`mlir-spirv-cpu-runner` is the name of a formerly separate tool that allows to
+execute `gpu` dialect kernel on the CPU via SPIR-V to LLVM dialect conversion.
+This runner is now merged with the `mlir-cpu-runner`; pass the
+`--link-nested-modules` flag. Currently, only single-threaded kernels are
supported.
-To build the runner, add the following option to `cmake`: `bash
--DMLIR_ENABLE_SPIRV_CPU_RUNNER=1`
+To build the required runtime libaries, add the following option to `cmake`:
+`-DMLIR_ENABLE_SPIRV_CPU_RUNNER=1`
### Pipeline
diff --git a/mlir/test/CMakeLists.txt b/mlir/test/CMakeLists.txt
index 361981605a76b7..d375e62fcc2c61 100644
--- a/mlir/test/CMakeLists.txt
+++ b/mlir/test/CMakeLists.txt
@@ -199,7 +199,6 @@ endif()
if(MLIR_ENABLE_SPIRV_CPU_RUNNER)
add_subdirectory(mlir-spirv-cpu-runner)
list(APPEND MLIR_TEST_DEPENDS
- mlir-spirv-cpu-runner
mlir_test_spirv_cpu_runner_c_wrappers
)
endif()
diff --git a/mlir/test/mlir-spirv-cpu-runner/double.mlir b/mlir/test/mlir-spirv-cpu-runner/double.mlir
index 35557ba1e94c00..6ea595a133ef5c 100644
--- a/mlir/test/mlir-spirv-cpu-runner/double.mlir
+++ b/mlir/test/mlir-spirv-cpu-runner/double.mlir
@@ -1,5 +1,5 @@
// RUN: mlir-opt %s -test-spirv-cpu-runner-pipeline \
-// RUN: | mlir-spirv-cpu-runner - -e main --entry-point-result=void --shared-libs=%mlir_runner_utils,%mlir_test_spirv_cpu_runner_c_wrappers \
+// RUN: | mlir-cpu-runner - -e main --entry-point-result=void --shared-libs=%mlir_runner_utils,%mlir_test_spirv_cpu_runner_c_wrappers --link-nested-modules \
// RUN: | FileCheck %s
// CHECK: [8, 8, 8, 8, 8, 8]
diff --git a/mlir/test/mlir-spirv-cpu-runner/simple_add.mlir b/mlir/test/mlir-spirv-cpu-runner/simple_add.mlir
index 75675a69a67583..f6de7fdb393935 100644
--- a/mlir/test/mlir-spirv-cpu-runner/simple_add.mlir
+++ b/mlir/test/mlir-spirv-cpu-runner/simple_add.mlir
@@ -1,5 +1,5 @@
// RUN: mlir-opt %s -test-spirv-cpu-runner-pipeline \
-// RUN: | mlir-spirv-cpu-runner - -e main --entry-point-result=void --shared-libs=%mlir_runner_utils,%mlir_test_spirv_cpu_runner_c_wrappers \
+// RUN: | mlir-cpu-runner - -e main --entry-point-result=void --shared-libs=%mlir_runner_utils,%mlir_test_spirv_cpu_runner_c_wrappers --link-nested-modules \
// RUN: | FileCheck %s
// CHECK: data =
diff --git a/mlir/tools/CMakeLists.txt b/mlir/tools/CMakeLists.txt
index 0a2d0ff2915099..072e83c5d45ea1 100644
--- a/mlir/tools/CMakeLists.txt
+++ b/mlir/tools/CMakeLists.txt
@@ -6,7 +6,6 @@ add_subdirectory(mlir-query)
add_subdirectory(mlir-reduce)
add_subdirectory(mlir-rewrite)
add_subdirectory(mlir-shlib)
-add_subdirectory(mlir-spirv-cpu-runner)
add_subdirectory(mlir-translate)
add_subdirectory(mlir-vulkan-runner)
add_subdirectory(tblgen-lsp-server)
diff --git a/mlir/tools/mlir-cpu-runner/mlir-cpu-runner.cpp b/mlir/tools/mlir-cpu-runner/mlir-cpu-runner.cpp
index 1b1d5d03d2be5d..b6486186b465f5 100644
--- a/mlir/tools/mlir-cpu-runner/mlir-cpu-runner.cpp
+++ b/mlir/tools/mlir-cpu-runner/mlir-cpu-runner.cpp
@@ -17,10 +17,65 @@
#include "mlir/ExecutionEngine/OptUtils.h"
#include "mlir/IR/Dialect.h"
#include "mlir/Target/LLVMIR/Dialect/All.h"
+#include "mlir/Target/LLVMIR/Dialect/LLVMIR/LLVMToLLVMIRTranslation.h"
+#include "mlir/Target/LLVMIR/Export.h"
+#include "llvm/IR/LLVMContext.h"
+#include "llvm/IR/Module.h"
+#include "llvm/Linker/Linker.h"
+#include "llvm/Support/CommandLine.h"
#include "llvm/Support/InitLLVM.h"
#include "llvm/Support/TargetSelect.h"
+using namespace mlir;
+
+llvm::cl::opt<bool> LinkNestedModules(
+ "link-nested-modules",
+ llvm::cl::desc("Link two nested MLIR modules into a single LLVM IR module. "
+ "Useful if both the host and device code can be run on the "
+ "same CPU, as in mlir-spirv-cpu-runner tests."));
+
+/// A utility function that builds llvm::Module from two nested MLIR modules.
+///
+/// module @main {
+/// module @kernel {
+/// // Some ops
+/// }
+/// // Some other ops
+/// }
+///
+/// Each of these two modules is translated to LLVM IR module, then they are
+/// linked together and returned.
+static std::unique_ptr<llvm::Module>
+convertMLIRModule(Operation *op, llvm::LLVMContext &context) {
+ auto module = dyn_cast<ModuleOp>(op);
+ if (!module)
+ return op->emitError("op must be a 'builtin.module"), nullptr;
+
+ std::unique_ptr<llvm::Module> kernelModule;
+ if (LinkNestedModules) {
+ // Verify that there is only one nested module.
+ auto modules = module.getOps<ModuleOp>();
+ if (!llvm::hasSingleElement(modules)) {
+ module.emitError("The module must contain exactly one nested module");
+ return nullptr;
+ }
+
+ // Translate nested module and erase it.
+ ModuleOp nested = *modules.begin();
+ kernelModule = translateModuleToLLVMIR(nested, context);
+ nested.erase();
+ }
+
+ std::unique_ptr<llvm::Module> mainModule =
+ translateModuleToLLVMIR(module, context);
+
+ if (LinkNestedModules)
+ llvm::Linker::linkModules(*mainModule, std::move(kernelModule));
+
+ return mainModule;
+}
+
int main(int argc, char **argv) {
llvm::InitLLVM y(argc, argv);
llvm::InitializeNativeTarget();
@@ -30,5 +85,7 @@ int main(int argc, char **argv) {
mlir::DialectRegistry registry;
mlir::registerAllToLLVMIRTranslations(registry);
- return mlir::JitRunnerMain(argc, argv, registry);
+ mlir::JitRunnerConfig jitRunnerConfig;
+ jitRunnerConfig.llvmModuleBuilder = convertMLIRModule;
+ return mlir::JitRunnerMain(argc, argv, registry, jitRunnerConfig);
}
diff --git a/mlir/tools/mlir-spirv-cpu-runner/CMakeLists.txt b/mlir/tools/mlir-spirv-cpu-runner/CMakeLists.txt
deleted file mode 100644
index 5760fad25ec65a..00000000000000
--- a/mlir/tools/mlir-spirv-cpu-runner/CMakeLists.txt
+++ /dev/null
@@ -1,36 +0,0 @@
-set(LLVM_LINK_COMPONENTS
- Linker
-)
-
-if (MLIR_ENABLE_SPIRV_CPU_RUNNER)
- message(STATUS "Building SPIR-V CPU runner")
-
- add_mlir_tool(mlir-spirv-cpu-runner
- mlir-spirv-cpu-runner.cpp
- )
-
- llvm_update_compile_flags(mlir-spirv-cpu-runner)
-
- get_property(conversion_libs GLOBAL PROPERTY MLIR_CONVERSION_LIBS)
-
- target_link_libraries(mlir-spirv-cpu-runner PRIVATE
- ${conversion_libs}
- MLIRAnalysis
- MLIRArithDialect
- MLIRBuiltinToLLVMIRTranslation
- MLIRExecutionEngine
- MLIRFuncDialect
- MLIRGPUDialect
- MLIRIR
- MLIRJitRunner
- MLIRLLVMDialect
- MLIRLLVMToLLVMIRTranslation
- MLIRMemRefDialect
- MLIRParser
- MLIRSPIRVDialect
- MLIRTargetLLVMIRExport
- MLIRTransforms
- MLIRTranslateLib
- MLIRSupport
- )
-endif()
diff --git a/mlir/tools/mlir-spirv-cpu-runner/mlir-spirv-cpu-runner.cpp b/mlir/tools/mlir-spirv-cpu-runner/mlir-spirv-cpu-runner.cpp
deleted file mode 100644
index 22ad1024db4a0b..00000000000000
--- a/mlir/tools/mlir-spirv-cpu-runner/mlir-spirv-cpu-runner.cpp
+++ /dev/null
@@ -1,90 +0,0 @@
-//===- mlir-spirv-cpu-runner.cpp - MLIR SPIR-V Execution on CPU -----------===//
-//
-// 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
-//
-//===----------------------------------------------------------------------===//
-//
-// Main entry point to a command line utility that executes an MLIR file on the
-// CPU by translating MLIR GPU module and host part to LLVM IR before
-// JIT-compiling and executing.
-//
-//===----------------------------------------------------------------------===//
-
-#include "mlir/Dialect/Arith/IR/Arith.h"
-#include "mlir/Dialect/Func/IR/FuncOps.h"
-#include "mlir/Dialect/GPU/IR/GPUDialect.h"
-#include "mlir/Dialect/LLVMIR/LLVMDialect.h"
-#include "mlir/Dialect/MemRef/IR/MemRef.h"
-#include "mlir/Dialect/SPIRV/IR/SPIRVDialect.h"
-#include "mlir/ExecutionEngine/JitRunner.h"
-#include "mlir/ExecutionEngine/OptUtils.h"
-#include "mlir/Pass/Pass.h"
-#include "mlir/Pass/PassManager.h"
-#include "mlir/Target/LLVMIR/Dialect/Builtin/BuiltinToLLVMIRTranslation.h"
-#include "mlir/Target/LLVMIR/Dialect/LLVMIR/LLVMToLLVMIRTranslation.h"
-#include "mlir/Target/LLVMIR/Export.h"
-
-#include "llvm/IR/LLVMContext.h"
-#include "llvm/IR/Module.h"
-#include "llvm/Linker/Linker.h"
-#include "llvm/Support/InitLLVM.h"
-#include "llvm/Support/TargetSelect.h"
-
-using namespace mlir;
-
-/// A utility function that builds llvm::Module from two nested MLIR modules.
-///
-/// module @main {
-/// module @kernel {
-/// // Some ops
-/// }
-/// // Some other ops
-/// }
-///
-/// Each of these two modules is translated to LLVM IR module, then they are
-/// linked together and returned.
-static std::unique_ptr<llvm::Module>
-convertMLIRModule(Operation *op, llvm::LLVMContext &context) {
- auto module = dyn_cast<ModuleOp>(op);
- if (!module)
- return op->emitError("op must be a 'builtin.module"), nullptr;
- // Verify that there is only one nested module.
- auto modules = module.getOps<ModuleOp>();
- if (!llvm::hasSingleElement(modules)) {
- module.emitError("The module must contain exactly one nested module");
- return nullptr;
- }
-
- // Translate nested module and erase it.
- ModuleOp nested = *modules.begin();
- std::unique_ptr<llvm::Module> kernelModule =
- translateModuleToLLVMIR(nested, context);
- nested.erase();
-
- std::unique_ptr<llvm::Module> mainModule =
- translateModuleToLLVMIR(module, context);
- llvm::Linker::linkModules(*mainModule, std::move(kernelModule));
- return mainModule;
-}
-
-int main(int argc, char **argv) {
- llvm::InitLLVM y(argc, argv);
-
- llvm::InitializeNativeTarget();
- llvm::InitializeNativeTargetAsmPrinter();
-
- mlir::JitRunnerConfig jitRunnerConfig;
- jitRunnerConfig.llvmModuleBuilder = convertMLIRModule;
-
- mlir::DialectRegistry registry;
- registry.insert<mlir::arith::ArithDialect, mlir::LLVM::LLVMDialect,
- mlir::gpu::GPUDialect, mlir::spirv::SPIRVDialect,
- mlir::func::FuncDialect, mlir::memref::MemRefDialect>();
- mlir::registerPassManagerCLOptions();
- mlir::registerBuiltinDialectTranslation(registry);
- mlir::registerLLVMDialectTranslation(registry);
-
- return mlir::JitRunnerMain(argc, argv, registry, jitRunnerConfig);
-}
``````````
</details>
https://github.com/llvm/llvm-project/pull/114563
More information about the Mlir-commits
mailing list