[Mlir-commits] [mlir] [mlir][mlir-spirv-cpu-runner] Move MLIR pass pipeline to mlir-opt (PR #111575)

llvmlistbot at llvm.org llvmlistbot at llvm.org
Tue Oct 8 11:26:23 PDT 2024 

llvmbot wrote:


<!--LLVM PR SUMMARY COMMENT-->

@llvm/pr-subscribers-mlir-spirv

Author: Andrea Faulds (andfau-amd)

<details>
<summary>Changes</summary>

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.

---
Full diff: https://github.com/llvm/llvm-project/pull/111575.diff


6 Files Affected:

- (modified) mlir/test/lib/Pass/CMakeLists.txt (+1) 
- (added) mlir/test/lib/Pass/TestSPIRVCPURunnerPipeline.cpp (+87) 
- (modified) mlir/test/mlir-spirv-cpu-runner/double.mlir (+2-1) 
- (modified) mlir/test/mlir-spirv-cpu-runner/simple_add.mlir (+2-1) 
- (modified) mlir/tools/mlir-opt/mlir-opt.cpp (+2) 
- (modified) mlir/tools/mlir-spirv-cpu-runner/mlir-spirv-cpu-runner.cpp (-24) 


``````````diff
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 &registry) 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;

``````````

</details>


https://github.com/llvm/llvm-project/pull/111575

More information about the Mlir-commits mailing list