[mlir] Remove the mlir-spirv-cpu-runner (move to mlir-cpu-runner) #114563
Conversation
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
|
@llvm/pr-subscribers-mlir-gpu @llvm/pr-subscribers-mlir-core Author: Andrea Faulds (andfau-amd) ChangesThis commit builds on and completes the work done in 9f6c632 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 ( Full diff: https://github.com/llvm/llvm-project/pull/114563.diff 8 Files Affected:
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);
-}
|
|
@llvm/pr-subscribers-mlir Author: Andrea Faulds (andfau-amd) ChangesThis commit builds on and completes the work done in 9f6c632 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 ( Full diff: https://github.com/llvm/llvm-project/pull/114563.diff 8 Files Affected:
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);
-}
|
|
Oh, should I move the "wrappers" to the execution engine? Is that just a renaming process? |
|
Nice! It's great to see the progress on this, thanks for pushing this forward :) |
Answering myself here: yes I should, as outlined in #73457, and yes it is, if https://reviews.llvm.org/D97463 is the clearest example of the pattern required. I do not have the clarity of mind needed right at this moment to execute it, but I should have this done tomorrow. I think it should be straightforward. |
andfau-amd
force-pushed
the
73457-runner-migration-cpu-spirv-remove
branch
from
2e7e1e9 to
4bc51c0
Compare
andfau-amd
force-pushed
the
73457-runner-migration-cpu-spirv-remove
branch
from
4bc51c0 to
c79a397
Compare
|
Okay, I finally got the renaming and moving done. It took me a bit to fully understand what the changes meant; I'm glad the CUDA example was so clear, because the other example I could find had too much stuff in one commit. I have also slightly reformatted and adjusted the CMake build definitions to match the style of the other runtimes. I would appreciate someone else testing this locally (I guess seeing if |
joker-eph
reviewed
Nov 6, 2024
kuhar
requested changes
Nov 6, 2024
There was a problem hiding this comment.
I agree with @joker-eph, the link stage should be done at the mlir-opt level, unless we have a very good reason for linking llvm IR over MLIR
|
Can we just call it |
I do think this would be a sensible change, but I think it should be done once all the separate runners are gone, as its own commit. |
kuhar
approved these changes
Nov 7, 2024
There was a problem hiding this comment.
LGTM % the documentation.
But let's also wait for a confirmation from @joker-eph.
andfau-amd
force-pushed
the
73457-runner-migration-cpu-spirv-remove
branch
from
c79a397 to
8e93963
Compare
This commit builds on and completes the work done in 9f6c632 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 remove the runner component of the mlir-spirv-cpu-runner. The runtime libraries and the tests are moved and renamed to fit into the Execution Engine and Integration tests, following the model of the similar migration done for the CUDA Runner in D97463.
andfau-amd
force-pushed
the
73457-runner-migration-cpu-spirv-remove
branch
from
8e93963 to
73e3b05
Compare
andfau-amd
commented
Nov 7, 2024
Comment on lines
+32
to
+34
| // TODO: Consider removing this linking functionality from the SPIR-V CPU Runner | ||
| // flow in favour of a more proper host/device split like other runners. | ||
| // https://github.com/llvm/llvm-project/issues/115348 |
There was a problem hiding this comment.
Here's the new TODO (#115348)
Groverkss
pushed a commit
to iree-org/llvm-project
that referenced
this pull request
Nov 15, 2024
…vm#114563) This commit builds on and completes the work done in 9f6c632 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 remove the runner component of the mlir-spirv-cpu-runner. The runtime libraries and the tests are moved and renamed to fit into the Execution Engine and Integration tests, following the model of the similar migration done for the CUDA Runner in D97463.
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.
This suggestion is invalid because no changes were made to the code.
Suggestions cannot be applied while the pull request is closed.
Suggestions cannot be applied while viewing a subset of changes.
Only one suggestion per line can be applied in a batch.
Add this suggestion to a batch that can be applied as a single commit.
Applying suggestions on deleted lines is not supported.
You must change the existing code in this line in order to create a valid suggestion.
Outdated suggestions cannot be applied.
This suggestion has been applied or marked resolved.
Suggestions cannot be applied from pending reviews.
Suggestions cannot be applied on multi-line comments.
Suggestions cannot be applied while the pull request is queued to merge.
Suggestion cannot be applied right now. Please check back later.
This commit builds on and completes the work done in 9f6c632 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 remove the runner component of the mlir-spirv-cpu-runner.The runtime libraries and the tests are moved and renamed to fit into the Execution Engine and Integration tests, following the model of the similar migration done for the CUDA Runner in D97463.