[Mlir-commits] [mlir] [mlir] Remove the mlir-spirv-cpu-runner (move to mlir-cpu-runner) (PR #114563)
Andrea Faulds
llvmlistbot at llvm.org
Thu Nov 7 09:22:35 PST 2024
https://github.com/andfau-amd updated https://github.com/llvm/llvm-project/pull/114563
>From 8e93963d3b4b11c988f179e6a02c58da0a0a1509 Mon Sep 17 00:00:00 2001
From: Andrea Faulds <andrea.faulds at amd.com>
Date: Thu, 7 Nov 2024 18:21:54 +0100
Subject: [PATCH] [mlir] Remove the mlir-spirv-cpu-runner (move to
mlir-cpu-runner)
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 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.
---
mlir/docs/SPIRVToLLVMDialectConversion.md | 17 ++--
mlir/lib/ExecutionEngine/CMakeLists.txt | 14 +++
.../SpirvCpuRuntimeWrappers.cpp} | 2 +-
mlir/test/CMakeLists.txt | 4 +-
.../GPU/SPIRV}/double.mlir | 2 +-
.../GPU/SPIRV}/lit.local.cfg | 0
.../GPU/SPIRV}/simple_add.mlir | 2 +-
.../lib/Pass/TestSPIRVCPURunnerPipeline.cpp | 4 +-
mlir/test/lit.cfg.py | 8 +-
.../test/mlir-spirv-cpu-runner/CMakeLists.txt | 6 --
mlir/tools/CMakeLists.txt | 1 -
.../tools/mlir-cpu-runner/mlir-cpu-runner.cpp | 59 +++++++++++-
.../mlir-spirv-cpu-runner/CMakeLists.txt | 36 --------
.../mlir-spirv-cpu-runner.cpp | 90 -------------------
14 files changed, 90 insertions(+), 155 deletions(-)
rename mlir/{test/mlir-spirv-cpu-runner/mlir_test_spirv_cpu_runner_c_wrappers.cpp => lib/ExecutionEngine/SpirvCpuRuntimeWrappers.cpp} (94%)
rename mlir/test/{mlir-spirv-cpu-runner => Integration/GPU/SPIRV}/double.mlir (94%)
rename mlir/test/{mlir-spirv-cpu-runner => Integration/GPU/SPIRV}/lit.local.cfg (100%)
rename mlir/test/{mlir-spirv-cpu-runner => Integration/GPU/SPIRV}/simple_add.mlir (95%)
delete mode 100644 mlir/test/mlir-spirv-cpu-runner/CMakeLists.txt
delete mode 100644 mlir/tools/mlir-spirv-cpu-runner/CMakeLists.txt
delete mode 100644 mlir/tools/mlir-spirv-cpu-runner/mlir-spirv-cpu-runner.cpp
diff --git a/mlir/docs/SPIRVToLLVMDialectConversion.md b/mlir/docs/SPIRVToLLVMDialectConversion.md
index 0aae02cff26be1..872690bb52646e 100644
--- a/mlir/docs/SPIRVToLLVMDialectConversion.md
+++ b/mlir/docs/SPIRVToLLVMDialectConversion.md
@@ -501,7 +501,7 @@ The SPIR-V to LLVM conversion does not involve modelling of workgroups. Hence,
we say that only current invocation is in conversion's scope. This means that
global variables with pointers of `Input`, `Output`, and `Private` storage
classes are supported. Also, `StorageBuffer` storage class is allowed for
-executing [`mlir-spirv-cpu-runner`](#mlir-spirv-cpu-runner).
+executing [SPIR-V CPU Runner tests](#spir-v-cpu-runner-tests).
Moreover, `bind` that specifies the descriptor set and the binding number and
`built_in` that specifies SPIR-V `BuiltIn` decoration have no conversion into
@@ -815,14 +815,15 @@ Module in SPIR-V has one region that contains one block. It is defined via
`spirv.module` is converted into `ModuleOp`. This plays a role of enclosing scope
to LLVM ops. At the moment, SPIR-V module attributes are ignored.
-## `mlir-spirv-cpu-runner`
+## SPIR-V CPU Runner Tests
-`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
-supported.
+The `mlir-cpu-runner` has support for executing a `gpu` dialect kernel on the
+CPU via SPIR-V to LLVM dialect conversion. This is referred to as the "SPIR-V
+CPU Runner". The `--link-nested-modules` flag needs to be passed for this.
+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
@@ -846,7 +847,7 @@ executed using `ExecutionEngine`.
### Walk-through
This section gives a detailed overview of the IR changes while running
-`mlir-spirv-cpu-runner`. First, consider that we have the following IR. (For
+SPIR-V CPU Runner tests. First, consider that we have the following IR. (For
simplicity some type annotations and function implementations have been
omitted).
diff --git a/mlir/lib/ExecutionEngine/CMakeLists.txt b/mlir/lib/ExecutionEngine/CMakeLists.txt
index 7fc17b97f0c56f..cf44a02cf5cb94 100644
--- a/mlir/lib/ExecutionEngine/CMakeLists.txt
+++ b/mlir/lib/ExecutionEngine/CMakeLists.txt
@@ -14,6 +14,7 @@ set(LLVM_OPTIONAL_SOURCES
RunnerUtils.cpp
OptUtils.cpp
JitRunner.cpp
+ SpirvCpuRuntimeWrappers.cpp
SyclRuntimeWrappers.cpp
)
@@ -401,4 +402,17 @@ if(LLVM_ENABLE_PIC)
set_property(TARGET mlir_sycl_runtime APPEND PROPERTY BUILD_RPATH "${LevelZero_LIBRARIES_DIR}" "${SyclRuntime_LIBRARIES_DIR}")
endif()
+
+ if(MLIR_ENABLE_SPIRV_CPU_RUNNER)
+ add_mlir_library(mlir_spirv_cpu_runtime
+ SHARED
+ SpirvCpuRuntimeWrappers.cpp
+
+ EXCLUDE_FROM_LIBMLIR
+ )
+
+ target_compile_definitions(mlir_spirv_cpu_runtime
+ PRIVATE
+ mlir_spirv_cpu_runtime_EXPORTS)
+ endif()
endif()
diff --git a/mlir/test/mlir-spirv-cpu-runner/mlir_test_spirv_cpu_runner_c_wrappers.cpp b/mlir/lib/ExecutionEngine/SpirvCpuRuntimeWrappers.cpp
similarity index 94%
rename from mlir/test/mlir-spirv-cpu-runner/mlir_test_spirv_cpu_runner_c_wrappers.cpp
rename to mlir/lib/ExecutionEngine/SpirvCpuRuntimeWrappers.cpp
index 93433be2dd66a6..e6594df3b1ee75 100644
--- a/mlir/test/mlir-spirv-cpu-runner/mlir_test_spirv_cpu_runner_c_wrappers.cpp
+++ b/mlir/lib/ExecutionEngine/SpirvCpuRuntimeWrappers.cpp
@@ -1,4 +1,4 @@
-//===- mlir_test_spirv_cpu_runner_c_wrappers.cpp - Runner testing library -===//
+//===- SpirvCpuRuntimeWrappers.cpp - Runner testing library -===//
//
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
// See https://llvm.org/LICENSE.txt for license information.
diff --git a/mlir/test/CMakeLists.txt b/mlir/test/CMakeLists.txt
index 361981605a76b7..f181a91328f3fe 100644
--- a/mlir/test/CMakeLists.txt
+++ b/mlir/test/CMakeLists.txt
@@ -197,10 +197,8 @@ if(LLVM_BUILD_EXAMPLES)
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
+ mlir_spirv_cpu_runtime
)
endif()
diff --git a/mlir/test/mlir-spirv-cpu-runner/double.mlir b/mlir/test/Integration/GPU/SPIRV/double.mlir
similarity index 94%
rename from mlir/test/mlir-spirv-cpu-runner/double.mlir
rename to mlir/test/Integration/GPU/SPIRV/double.mlir
index 35557ba1e94c00..b6887c26ebe0aa 100644
--- a/mlir/test/mlir-spirv-cpu-runner/double.mlir
+++ b/mlir/test/Integration/GPU/SPIRV/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_spirv_cpu_runtime --link-nested-modules \
// RUN: | FileCheck %s
// CHECK: [8, 8, 8, 8, 8, 8]
diff --git a/mlir/test/mlir-spirv-cpu-runner/lit.local.cfg b/mlir/test/Integration/GPU/SPIRV/lit.local.cfg
similarity index 100%
rename from mlir/test/mlir-spirv-cpu-runner/lit.local.cfg
rename to mlir/test/Integration/GPU/SPIRV/lit.local.cfg
diff --git a/mlir/test/mlir-spirv-cpu-runner/simple_add.mlir b/mlir/test/Integration/GPU/SPIRV/simple_add.mlir
similarity index 95%
rename from mlir/test/mlir-spirv-cpu-runner/simple_add.mlir
rename to mlir/test/Integration/GPU/SPIRV/simple_add.mlir
index 75675a69a67583..0049ce60c4ebbc 100644
--- a/mlir/test/mlir-spirv-cpu-runner/simple_add.mlir
+++ b/mlir/test/Integration/GPU/SPIRV/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_spirv_cpu_runtime --link-nested-modules \
// RUN: | FileCheck %s
// CHECK: data =
diff --git a/mlir/test/lib/Pass/TestSPIRVCPURunnerPipeline.cpp b/mlir/test/lib/Pass/TestSPIRVCPURunnerPipeline.cpp
index ded0d22c31307e..bf2880aa022eed 100644
--- a/mlir/test/lib/Pass/TestSPIRVCPURunnerPipeline.cpp
+++ b/mlir/test/lib/Pass/TestSPIRVCPURunnerPipeline.cpp
@@ -6,7 +6,7 @@
//
//===----------------------------------------------------------------------===//
//
-// Implements a pipeline for use by mlir-spirv-cpu-runner tests.
+// Implements a pipeline for use by SPIR-V CPU Runner tests.
//
//===----------------------------------------------------------------------===//
@@ -40,7 +40,7 @@ void registerTestSPIRVCPURunnerPipeline() {
PassPipelineRegistration<>(
"test-spirv-cpu-runner-pipeline",
"Runs a series of passes for lowering SPIR-V-dialect MLIR to "
- "LLVM-dialect MLIR intended for mlir-spirv-cpu-runner.",
+ "LLVM-dialect MLIR intended for SPIR-V CPU Runner tests.",
buildTestSPIRVCPURunnerPipeline);
}
} // namespace test
diff --git a/mlir/test/lit.cfg.py b/mlir/test/lit.cfg.py
index f0d4f35ba3e229..9b429b424d3575 100644
--- a/mlir/test/lit.cfg.py
+++ b/mlir/test/lit.cfg.py
@@ -125,11 +125,6 @@ def add_runtime(name):
"not",
]
-if config.enable_spirv_cpu_runner:
- tools.extend(
- ["mlir-spirv-cpu-runner", add_runtime("mlir_test_spirv_cpu_runner_c_wrappers")]
- )
-
if config.enable_vulkan_runner:
tools.extend([add_runtime("vulkan-runtime-wrappers")])
@@ -142,6 +137,9 @@ def add_runtime(name):
if config.enable_sycl_runner:
tools.extend([add_runtime("mlir_sycl_runtime")])
+if config.enable_spirv_cpu_runner:
+ tools.extend([add_runtime("mlir_spirv_cpu_runtime")])
+
if config.mlir_run_arm_sve_tests or config.mlir_run_arm_sme_tests:
tools.extend([add_runtime("mlir_arm_runner_utils")])
diff --git a/mlir/test/mlir-spirv-cpu-runner/CMakeLists.txt b/mlir/test/mlir-spirv-cpu-runner/CMakeLists.txt
deleted file mode 100644
index c7521db57a38e2..00000000000000
--- a/mlir/test/mlir-spirv-cpu-runner/CMakeLists.txt
+++ /dev/null
@@ -1,6 +0,0 @@
-set(LLVM_OPTIONAL_SOURCES
- mlir_test_spirv_cpu_runner_c_wrappers.cpp
- )
-
-add_llvm_library(mlir_test_spirv_cpu_runner_c_wrappers SHARED mlir_test_spirv_cpu_runner_c_wrappers.cpp)
-target_compile_definitions(mlir_test_spirv_cpu_runner_c_wrappers PRIVATE mlir_test_spirv_cpu_runner_c_wrappers_EXPORTS)
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..542239d4da97a1 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 SPIR-V 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);
-}
More information about the Mlir-commits
mailing list