-
Notifications
You must be signed in to change notification settings - Fork 13.6k
[mlir][mlir-spirv-cpu-runner] Move MLIR pass pipeline to mlir-opt #111575
New issue
Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.
By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.
Already on GitHub? Sign in to your account
[mlir][mlir-spirv-cpu-runner] Move MLIR pass pipeline to mlir-opt #111575
Conversation
@llvm/pr-subscribers-mlir-core @llvm/pr-subscribers-mlir-spirv Author: Andrea Faulds (andfau-amd) ChangesAdds a new mlir-opt test-only pass, -test-spirv-cpu-runner-pipeline, which runs the set of MLIR passes needed for the mlir-spirv-cpu-runner, and removes them from the runner. The tests are changed to invoke mlir-opt with this flag before running the runner. The eventual goal is to move all host/device code generation steps out of the runner, like with some of the other runners. Full diff: https://github.com/llvm/llvm-project/pull/111575.diff 6 Files Affected:
diff --git a/mlir/test/lib/Pass/CMakeLists.txt b/mlir/test/lib/Pass/CMakeLists.txt
index dd90c228cdaf5f..9f79944ff89684 100644
--- a/mlir/test/lib/Pass/CMakeLists.txt
+++ b/mlir/test/lib/Pass/CMakeLists.txt
@@ -2,6 +2,7 @@
add_mlir_library(MLIRTestPass
TestDynamicPipeline.cpp
TestPassManager.cpp
+ TestSPIRVCPURunnerPipeline.cpp
EXCLUDE_FROM_LIBMLIR
diff --git a/mlir/test/lib/Pass/TestSPIRVCPURunnerPipeline.cpp b/mlir/test/lib/Pass/TestSPIRVCPURunnerPipeline.cpp
new file mode 100644
index 00000000000000..514b6fa6e3fd2b
--- /dev/null
+++ b/mlir/test/lib/Pass/TestSPIRVCPURunnerPipeline.cpp
@@ -0,0 +1,87 @@
+//===------------------ TestSPIRVCPURunnerPipeline.cpp --------------------===//
+//
+// 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 for use by mlir-spirv-cpu-runner tests.
+//
+//===----------------------------------------------------------------------===//
+
+#include "mlir/Conversion/FuncToLLVM/ConvertFuncToLLVMPass.h"
+#include "mlir/Conversion/GPUToSPIRV/GPUToSPIRVPass.h"
+#include "mlir/Conversion/SPIRVToLLVM/SPIRVToLLVMPass.h"
+#include "mlir/Dialect/Arith/IR/Arith.h"
+#include "mlir/Dialect/DLTI/DLTI.h"
+#include "mlir/Dialect/Func/IR/FuncOps.h"
+#include "mlir/Dialect/GPU/IR/GPUDialect.h"
+#include "mlir/Dialect/GPU/Transforms/Passes.h"
+#include "mlir/Dialect/LLVMIR/LLVMDialect.h"
+#include "mlir/Dialect/MemRef/IR/MemRef.h"
+#include "mlir/Dialect/SPIRV/IR/SPIRVDialect.h"
+#include "mlir/Dialect/SPIRV/IR/SPIRVOps.h"
+#include "mlir/Dialect/SPIRV/Transforms/Passes.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"
+
+using namespace mlir;
+
+namespace {
+
+class TestSPIRVCPURunnerPipelinePass
+ : public PassWrapper<TestSPIRVCPURunnerPipelinePass,
+ OperationPass<ModuleOp>> {
+public:
+ MLIR_DEFINE_EXPLICIT_INTERNAL_INLINE_TYPE_ID(TestSPIRVCPURunnerPipelinePass)
+
+ StringRef getArgument() const final {
+ return "test-spirv-cpu-runner-pipeline";
+ }
+ StringRef getDescription() const final {
+ return "Runs a series of passes for lowering SPIR-V-dialect MLIR to "
+ "LLVM-dialect MLIR intended for mlir-spirv-cpu-runner.";
+ }
+ void getDependentDialects(DialectRegistry ®istry) const override {
+ registry.insert<mlir::arith::ArithDialect, mlir::LLVM::LLVMDialect,
+ mlir::gpu::GPUDialect, mlir::spirv::SPIRVDialect,
+ mlir::func::FuncDialect, mlir::memref::MemRefDialect,
+ mlir::DLTIDialect>();
+ }
+
+ TestSPIRVCPURunnerPipelinePass() = default;
+ TestSPIRVCPURunnerPipelinePass(const TestSPIRVCPURunnerPipelinePass &) {}
+
+ void runOnOperation() override {
+ ModuleOp module = getOperation();
+
+ PassManager passManager(module->getContext(),
+ module->getName().getStringRef());
+ if (failed(applyPassManagerCLOptions(passManager)))
+ return signalPassFailure();
+ passManager.addPass(createGpuKernelOutliningPass());
+ passManager.addPass(createConvertGPUToSPIRVPass(/*mapMemorySpace=*/true));
+
+ OpPassManager &nestedPM = passManager.nest<spirv::ModuleOp>();
+ nestedPM.addPass(spirv::createSPIRVLowerABIAttributesPass());
+ nestedPM.addPass(spirv::createSPIRVUpdateVCEPass());
+ passManager.addPass(createLowerHostCodeToLLVMPass());
+ passManager.addPass(createConvertSPIRVToLLVMPass());
+
+ if (failed(runPipeline(passManager, module)))
+ signalPassFailure();
+ }
+};
+} // namespace
+
+namespace mlir {
+namespace test {
+void registerTestSPIRVCPURunnerPipelinePass() {
+ PassRegistration<TestSPIRVCPURunnerPipelinePass>();
+}
+} // namespace test
+} // namespace mlir
diff --git a/mlir/test/mlir-spirv-cpu-runner/double.mlir b/mlir/test/mlir-spirv-cpu-runner/double.mlir
index cd551ffb1bd062..35557ba1e94c00 100644
--- a/mlir/test/mlir-spirv-cpu-runner/double.mlir
+++ b/mlir/test/mlir-spirv-cpu-runner/double.mlir
@@ -1,4 +1,5 @@
-// RUN: mlir-spirv-cpu-runner %s -e main --entry-point-result=void --shared-libs=%mlir_runner_utils,%mlir_test_spirv_cpu_runner_c_wrappers \
+// 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: | 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 119e973e45e4a7..75675a69a67583 100644
--- a/mlir/test/mlir-spirv-cpu-runner/simple_add.mlir
+++ b/mlir/test/mlir-spirv-cpu-runner/simple_add.mlir
@@ -1,4 +1,5 @@
-// RUN: mlir-spirv-cpu-runner %s -e main --entry-point-result=void --shared-libs=%mlir_runner_utils,%mlir_test_spirv_cpu_runner_c_wrappers \
+// 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: | FileCheck %s
// CHECK: data =
diff --git a/mlir/tools/mlir-opt/mlir-opt.cpp b/mlir/tools/mlir-opt/mlir-opt.cpp
index 36b142484bb04a..6b639829aacd2c 100644
--- a/mlir/tools/mlir-opt/mlir-opt.cpp
+++ b/mlir/tools/mlir-opt/mlir-opt.cpp
@@ -142,6 +142,7 @@ void registerTestSCFWhileOpBuilderPass();
void registerTestSCFWrapInZeroTripCheckPasses();
void registerTestShapeMappingPass();
void registerTestSliceAnalysisPass();
+void registerTestSPIRVCPURunnerPipelinePass();
void registerTestSPIRVFuncSignatureConversion();
void registerTestSPIRVVectorUnrolling();
void registerTestTensorCopyInsertionPass();
@@ -278,6 +279,7 @@ void registerTestPasses() {
mlir::test::registerTestSCFWrapInZeroTripCheckPasses();
mlir::test::registerTestShapeMappingPass();
mlir::test::registerTestSliceAnalysisPass();
+ mlir::test::registerTestSPIRVCPURunnerPipelinePass();
mlir::test::registerTestSPIRVFuncSignatureConversion();
mlir::test::registerTestSPIRVVectorUnrolling();
mlir::test::registerTestTensorCopyInsertionPass();
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
index 7e0b51cac80621..22ad1024db4a0b 100644
--- a/mlir/tools/mlir-spirv-cpu-runner/mlir-spirv-cpu-runner.cpp
+++ b/mlir/tools/mlir-spirv-cpu-runner/mlir-spirv-cpu-runner.cpp
@@ -12,18 +12,12 @@
//
//===----------------------------------------------------------------------===//
-#include "mlir/Conversion/FuncToLLVM/ConvertFuncToLLVMPass.h"
-#include "mlir/Conversion/GPUToSPIRV/GPUToSPIRVPass.h"
-#include "mlir/Conversion/SPIRVToLLVM/SPIRVToLLVMPass.h"
#include "mlir/Dialect/Arith/IR/Arith.h"
#include "mlir/Dialect/Func/IR/FuncOps.h"
#include "mlir/Dialect/GPU/IR/GPUDialect.h"
-#include "mlir/Dialect/GPU/Transforms/Passes.h"
#include "mlir/Dialect/LLVMIR/LLVMDialect.h"
#include "mlir/Dialect/MemRef/IR/MemRef.h"
#include "mlir/Dialect/SPIRV/IR/SPIRVDialect.h"
-#include "mlir/Dialect/SPIRV/IR/SPIRVOps.h"
-#include "mlir/Dialect/SPIRV/Transforms/Passes.h"
#include "mlir/ExecutionEngine/JitRunner.h"
#include "mlir/ExecutionEngine/OptUtils.h"
#include "mlir/Pass/Pass.h"
@@ -75,23 +69,6 @@ convertMLIRModule(Operation *op, llvm::LLVMContext &context) {
return mainModule;
}
-static LogicalResult runMLIRPasses(Operation *module,
- JitRunnerOptions &options) {
- PassManager passManager(module->getContext(),
- module->getName().getStringRef());
- if (failed(applyPassManagerCLOptions(passManager)))
- return failure();
- passManager.addPass(createGpuKernelOutliningPass());
- passManager.addPass(createConvertGPUToSPIRVPass(/*mapMemorySpace=*/true));
-
- OpPassManager &nestedPM = passManager.nest<spirv::ModuleOp>();
- nestedPM.addPass(spirv::createSPIRVLowerABIAttributesPass());
- nestedPM.addPass(spirv::createSPIRVUpdateVCEPass());
- passManager.addPass(createLowerHostCodeToLLVMPass());
- passManager.addPass(createConvertSPIRVToLLVMPass());
- return passManager.run(module);
-}
-
int main(int argc, char **argv) {
llvm::InitLLVM y(argc, argv);
@@ -99,7 +76,6 @@ int main(int argc, char **argv) {
llvm::InitializeNativeTargetAsmPrinter();
mlir::JitRunnerConfig jitRunnerConfig;
- jitRunnerConfig.mlirTransformer = runMLIRPasses;
jitRunnerConfig.llvmModuleBuilder = convertMLIRModule;
mlir::DialectRegistry registry;
|
For broader motivation see #73457. |
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
Nice!
7d5f2d2
to
cee317a
Compare
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
LGTM % nit
cee317a
to
1fcb1a9
Compare
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
Thanks for the fixes
9a445c6
to
0b2fdef
Compare
Adds a new mlir-opt test-only pipeline, -test-spirv-cpu-runner-pipeline, containing the set of MLIR passes needed for the mlir-spirv-cpu-runner, and removes them from the runner. The tests are changed to invoke mlir-opt with this flag before running the runner. The eventual goal is to move all host/device code generation steps out of the runner, like with some of the other runners.
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
LGTM
Thanks for the suggestion @joker-eph |
LLVM Buildbot has detected a new failure on builder Full details are available at: https://lab.llvm.org/buildbot/#/builders/89/builds/8820 Here is the relevant piece of the build log for the reference
|
LLVM Buildbot has detected a new failure on builder Full details are available at: https://lab.llvm.org/buildbot/#/builders/138/builds/5244 Here is the relevant piece of the build log for the reference
|
LLVM Buildbot has detected a new failure on builder Full details are available at: https://lab.llvm.org/buildbot/#/builders/80/builds/5191 Here is the relevant piece of the build log for the reference
|
LLVM Buildbot has detected a new failure on builder Full details are available at: https://lab.llvm.org/buildbot/#/builders/130/builds/5123 Here is the relevant piece of the build log for the reference
|
…-opt" (#113176) Reverts #111575 This caused build failures: https://lab.llvm.org/buildbot/#/builders/138/builds/5244
Adds a new mlir-opt test-only pass, -test-spirv-cpu-runner-pipeline, which runs the set of MLIR passes needed for the mlir-spirv-cpu-runner, and removes them from the runner. The tests are changed to invoke mlir-opt with this flag before running the runner. The eventual goal is to move all host/device code generation steps out of the runner, like with some of the other runners.