[mlir][spirv] Do SPIR-V serialization in -test-vulkan-runner-pipeline #121494
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 Author: Andrea Faulds (andfau-amd) ChangesThis commit is a further incremental step toward moving the whole mlir-vulkan-runner MLIR pass pipeline into mlir-opt (see #73457). The previous step was b225b3adf7b78387c9fcb97a3ff0e0a1e26eafe2, which moved all device passes prior to SPIR-V serialization into a new mlir-opt test pass, This commit changes how SPIR-V serialization is accomplished for Vulkan runner tests. Until now, this was done by the Vulkan-specific ConvertGpuLaunchFuncToVulkanLaunchFunc pass. With this commit, this responsibility is removed from that pass, and is instead done with the existing generic GpuModuleToBinaryPass. In addition, the SPIR-V serialization step is no longer done inside mlir-vulkan-runner, but rather inside mlir-opt (in the In order to enable this, this commit includes these supporting changes:
Full diff: https://github.com/llvm/llvm-project/pull/121494.diff 10 Files Affected:
diff --git a/mlir/include/mlir/Conversion/Passes.td b/mlir/include/mlir/Conversion/Passes.td
index 8835e0a9099fdd..e8713230814d01 100644
--- a/mlir/include/mlir/Conversion/Passes.td
+++ b/mlir/include/mlir/Conversion/Passes.td
@@ -61,7 +61,10 @@ def ConvertToSPIRVPass : Pass<"convert-to-spirv"> {
"Run vector unrolling to convert vector types in function bodies">,
Option<"convertGPUModules", "convert-gpu-modules", "bool",
/*default=*/"false",
- "Clone and convert GPU modules">
+ "Clone and convert GPU modules">,
+ Option<"nestInGPUModule", "nest-in-gpu-module", "bool",
+ /*default=*/"false",
+ "Put converted SPIR-V module inside the gpu.module instead of alongside it.">,
];
}
diff --git a/mlir/lib/Conversion/ConvertToSPIRV/ConvertToSPIRVPass.cpp b/mlir/lib/Conversion/ConvertToSPIRV/ConvertToSPIRVPass.cpp
index 4b7f7ff114deeb..ab9c048f561069 100644
--- a/mlir/lib/Conversion/ConvertToSPIRV/ConvertToSPIRVPass.cpp
+++ b/mlir/lib/Conversion/ConvertToSPIRV/ConvertToSPIRVPass.cpp
@@ -108,7 +108,10 @@ struct ConvertToSPIRVPass final
SmallVector<Operation *, 1> gpuModules;
OpBuilder builder(context);
op->walk([&](gpu::GPUModuleOp gpuModule) {
- builder.setInsertionPoint(gpuModule);
+ if (nestInGPUModule)
+ builder.setInsertionPointToStart(gpuModule.getBody());
+ else
+ builder.setInsertionPoint(gpuModule);
gpuModules.push_back(builder.clone(*gpuModule));
});
// Run conversion for each module independently as they can have
diff --git a/mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRVPass.cpp b/mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRVPass.cpp
index 08b451f7d5b325..509b6343057b99 100644
--- a/mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRVPass.cpp
+++ b/mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRVPass.cpp
@@ -71,7 +71,7 @@ void GPUToSPIRVPass::runOnOperation() {
// launch op still needs the original GPU kernel module.
// For Vulkan Shader capabilities, we insert the newly converted SPIR-V
// module right after the original GPU module, as that's the expectation of
- // the in-tree Vulkan runner.
+ // the in-tree SPIR-V CPU runner (the Vulkan runner does not use this pass).
// For OpenCL Kernel capabilities, we insert the newly converted SPIR-V
// module inside the original GPU module, as that's the expectaion of the
// normal GPU compilation pipeline.
diff --git a/mlir/lib/Conversion/GPUToVulkan/ConvertGPULaunchFuncToVulkanLaunchFunc.cpp b/mlir/lib/Conversion/GPUToVulkan/ConvertGPULaunchFuncToVulkanLaunchFunc.cpp
index 2d2251672230b6..69945cb6db8226 100644
--- a/mlir/lib/Conversion/GPUToVulkan/ConvertGPULaunchFuncToVulkanLaunchFunc.cpp
+++ b/mlir/lib/Conversion/GPUToVulkan/ConvertGPULaunchFuncToVulkanLaunchFunc.cpp
@@ -7,9 +7,8 @@
//===----------------------------------------------------------------------===//
//
// This file implements a pass to convert gpu launch function into a vulkan
-// launch function. Creates a SPIR-V binary shader from the `spirv::ModuleOp`
-// using `spirv::serialize` function, attaches binary data and entry point name
-// as an attributes to vulkan launch call op.
+// launch function. Extracts the SPIR-V from a `gpu::BinaryOp` and attaches it
+// along with the entry point name as attributes to a Vulkan launch call op.
//
//===----------------------------------------------------------------------===//
@@ -40,10 +39,9 @@ static constexpr const char *kVulkanLaunch = "vulkanLaunch";
namespace {
-/// A pass to convert gpu launch op to vulkan launch call op, by creating a
-/// SPIR-V binary shader from `spirv::ModuleOp` using `spirv::serialize`
-/// function and attaching binary data and entry point name as an attributes to
-/// created vulkan launch call op.
+/// A pass to convert gpu launch op to vulkan launch call op, by extracting a
+/// SPIR-V binary shader from a `gpu::BinaryOp` and attaching binary data and
+/// entry point name as an attributes to created vulkan launch call op.
class ConvertGpuLaunchFuncToVulkanLaunchFunc
: public impl::ConvertGpuLaunchFuncToVulkanLaunchFuncBase<
ConvertGpuLaunchFuncToVulkanLaunchFunc> {
@@ -51,10 +49,10 @@ class ConvertGpuLaunchFuncToVulkanLaunchFunc
void runOnOperation() override;
private:
- /// Creates a SPIR-V binary shader from the given `module` using
- /// `spirv::serialize` function.
- LogicalResult createBinaryShader(ModuleOp module,
- std::vector<char> &binaryShader);
+ /// Extracts a SPIR-V binary shader from the given `module`, if any.
+ /// Note that this also removes the binary from the IR.
+ LogicalResult getBinaryShader(ModuleOp module,
+ std::vector<char> &binaryShader);
/// Converts the given `launchOp` to vulkan launch call.
void convertGpuLaunchFunc(gpu::LaunchFuncOp launchOp);
@@ -135,21 +133,38 @@ LogicalResult ConvertGpuLaunchFuncToVulkanLaunchFunc::declareVulkanLaunchFunc(
return success();
}
-LogicalResult ConvertGpuLaunchFuncToVulkanLaunchFunc::createBinaryShader(
+LogicalResult ConvertGpuLaunchFuncToVulkanLaunchFunc::getBinaryShader(
ModuleOp module, std::vector<char> &binaryShader) {
bool done = false;
SmallVector<uint32_t, 0> binary;
- for (auto spirvModule : module.getOps<spirv::ModuleOp>()) {
+ gpu::BinaryOp *binaryToErase;
+ for (auto gpuBinary : module.getOps<gpu::BinaryOp>()) {
if (done)
- return spirvModule.emitError("should only contain one 'spirv.module' op");
+ return gpuBinary.emitError("should only contain one 'gpu.binary' op");
done = true;
- if (failed(spirv::serialize(spirvModule, binary)))
- return failure();
+ ArrayRef<Attribute> objects = gpuBinary.getObjectsAttr().getValue();
+ if (objects.size() != 1)
+ return gpuBinary.emitError("should only contain a single object");
+
+ auto object = cast<gpu::ObjectAttr>(objects[0]);
+
+ if (!isa<spirv::TargetEnvAttr>(object.getTarget()))
+ return gpuBinary.emitError(
+ "should contain an object with a SPIR-V target environment");
+
+ StringAttr objectStrAttr = object.getObject();
+ StringRef objectStr = objectStrAttr.getValue();
+ binaryShader.insert(binaryShader.end(), objectStr.bytes_begin(),
+ objectStr.bytes_end());
+
+ binaryToErase = &gpuBinary;
}
- binaryShader.resize(binary.size() * sizeof(uint32_t));
- std::memcpy(binaryShader.data(), reinterpret_cast<char *>(binary.data()),
- binaryShader.size());
+ if (!done)
+ return module.emitError("should contain a 'gpu.binary' op");
+
+ // Remove the binary to avoid confusing later conversion passes.
+ binaryToErase->erase();
return success();
}
@@ -159,9 +174,9 @@ void ConvertGpuLaunchFuncToVulkanLaunchFunc::convertGpuLaunchFunc(
OpBuilder builder(launchOp);
Location loc = launchOp.getLoc();
- // Serialize `spirv::Module` into binary form.
std::vector<char> binary;
- if (failed(createBinaryShader(module, binary)))
+ // Extract SPIR-V from `gpu.binary` op.
+ if (failed(getBinaryShader(module, binary)))
return signalPassFailure();
// Declare vulkan launch function.
diff --git a/mlir/test/Conversion/ConvertToSPIRV/convert-gpu-modules-nested.mlir b/mlir/test/Conversion/ConvertToSPIRV/convert-gpu-modules-nested.mlir
new file mode 100644
index 00000000000000..33fa0f859a5c77
--- /dev/null
+++ b/mlir/test/Conversion/ConvertToSPIRV/convert-gpu-modules-nested.mlir
@@ -0,0 +1,30 @@
+// RUN: mlir-opt -convert-to-spirv="convert-gpu-modules=true nest-in-gpu-module=true run-signature-conversion=false run-vector-unrolling=false" -split-input-file %s | FileCheck %s
+
+module attributes {
+ gpu.container_module,
+ spirv.target_env = #spirv.target_env<#spirv.vce<v1.0, [Shader], []>, #spirv.resource_limits<>>
+} {
+ // CHECK-LABEL: func.func @main
+ // CHECK: %[[C1:.*]] = arith.constant 1 : index
+ // CHECK: gpu.launch_func @[[$KERNELS_1:.*]]::@[[$BUILTIN_WG_ID_X:.*]] blocks in (%[[C1]], %[[C1]], %[[C1]]) threads in (%[[C1]], %[[C1]], %[[C1]])
+ func.func @main() {
+ %c1 = arith.constant 1 : index
+ gpu.launch_func @kernels_1::@builtin_workgroup_id_x
+ blocks in (%c1, %c1, %c1) threads in (%c1, %c1, %c1)
+ return
+ }
+
+ // CHECK: gpu.module @[[$KERNELS_1]]
+ // CHECK: spirv.module @{{.*}} Logical GLSL450
+ // CHECK: spirv.func @[[$BUILTIN_WG_ID_X]]
+ // CHECK: spirv.mlir.addressof
+ // CHECK: spirv.Load "Input"
+ // CHECK: spirv.CompositeExtract
+ gpu.module @kernels_1 {
+ gpu.func @builtin_workgroup_id_x() kernel
+ attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {
+ %0 = gpu.block_id x
+ gpu.return
+ }
+ }
+}
diff --git a/mlir/test/Conversion/GPUToVulkan/lower-gpu-launch-vulkan-launch.mlir b/mlir/test/Conversion/GPUToVulkan/lower-gpu-launch-vulkan-launch.mlir
index 665d0a33abedc2..96ee1866517e6d 100644
--- a/mlir/test/Conversion/GPUToVulkan/lower-gpu-launch-vulkan-launch.mlir
+++ b/mlir/test/Conversion/GPUToVulkan/lower-gpu-launch-vulkan-launch.mlir
@@ -1,24 +1,24 @@
-// RUN: mlir-opt %s -convert-gpu-launch-to-vulkan-launch | FileCheck %s
+// RUN: mlir-opt %s -pass-pipeline='builtin.module(spirv-attach-target{ver=v1.0 caps=Shader exts=SPV_KHR_storage_buffer_storage_class},gpu-module-to-binary,convert-gpu-launch-to-vulkan-launch)' | FileCheck %s
// CHECK: %[[resource:.*]] = memref.alloc() : memref<12xf32>
// CHECK: %[[index:.*]] = arith.constant 1 : index
// CHECK: call @vulkanLaunch(%[[index]], %[[index]], %[[index]], %[[resource]]) {spirv_blob = "{{.*}}", spirv_element_types = [f32], spirv_entry_point = "kernel"}
module attributes {gpu.container_module} {
- spirv.module Logical GLSL450 requires #spirv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]> {
- spirv.GlobalVariable @kernel_arg_0 bind(0, 0) : !spirv.ptr<!spirv.struct<(!spirv.array<12 x f32, stride=4> [0])>, StorageBuffer>
- spirv.func @kernel() "None" attributes {workgroup_attributions = 0 : i64} {
- %0 = spirv.mlir.addressof @kernel_arg_0 : !spirv.ptr<!spirv.struct<(!spirv.array<12 x f32, stride=4> [0])>, StorageBuffer>
- %2 = spirv.Constant 0 : i32
- %3 = spirv.mlir.addressof @kernel_arg_0 : !spirv.ptr<!spirv.struct<(!spirv.array<12 x f32, stride=4> [0])>, StorageBuffer>
- %4 = spirv.AccessChain %0[%2, %2] : !spirv.ptr<!spirv.struct<(!spirv.array<12 x f32, stride=4> [0])>, StorageBuffer>, i32, i32 -> !spirv.ptr<f32, StorageBuffer>
- %5 = spirv.Load "StorageBuffer" %4 : f32
- spirv.Return
- }
- spirv.EntryPoint "GLCompute" @kernel
- spirv.ExecutionMode @kernel "LocalSize", 1, 1, 1
- }
gpu.module @kernels {
+ spirv.module Logical GLSL450 requires #spirv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]> {
+ spirv.GlobalVariable @kernel_arg_0 bind(0, 0) : !spirv.ptr<!spirv.struct<(!spirv.array<12 x f32, stride=4> [0])>, StorageBuffer>
+ spirv.func @kernel() "None" attributes {workgroup_attributions = 0 : i64} {
+ %0 = spirv.mlir.addressof @kernel_arg_0 : !spirv.ptr<!spirv.struct<(!spirv.array<12 x f32, stride=4> [0])>, StorageBuffer>
+ %2 = spirv.Constant 0 : i32
+ %3 = spirv.mlir.addressof @kernel_arg_0 : !spirv.ptr<!spirv.struct<(!spirv.array<12 x f32, stride=4> [0])>, StorageBuffer>
+ %4 = spirv.AccessChain %0[%2, %2] : !spirv.ptr<!spirv.struct<(!spirv.array<12 x f32, stride=4> [0])>, StorageBuffer>, i32, i32 -> !spirv.ptr<f32, StorageBuffer>
+ %5 = spirv.Load "StorageBuffer" %4 : f32
+ spirv.Return
+ }
+ spirv.EntryPoint "GLCompute" @kernel
+ spirv.ExecutionMode @kernel "LocalSize", 1, 1, 1
+ }
gpu.func @kernel(%arg0: memref<12xf32>) kernel {
gpu.return
}
diff --git a/mlir/test/lib/Pass/TestVulkanRunnerPipeline.cpp b/mlir/test/lib/Pass/TestVulkanRunnerPipeline.cpp
index eda9aa9f9efef7..9bd4c42a1cdfb5 100644
--- a/mlir/test/lib/Pass/TestVulkanRunnerPipeline.cpp
+++ b/mlir/test/lib/Pass/TestVulkanRunnerPipeline.cpp
@@ -12,33 +12,57 @@
#include "mlir/Conversion/ConvertToSPIRV/ConvertToSPIRVPass.h"
#include "mlir/Conversion/GPUToSPIRV/GPUToSPIRVPass.h"
+#include "mlir/Dialect/GPU/IR/GPUDialect.h"
#include "mlir/Dialect/GPU/Transforms/Passes.h"
#include "mlir/Dialect/MemRef/Transforms/Passes.h"
#include "mlir/Dialect/SPIRV/IR/SPIRVOps.h"
#include "mlir/Dialect/SPIRV/Transforms/Passes.h"
#include "mlir/Pass/PassManager.h"
+#include "mlir/Pass/PassOptions.h"
using namespace mlir;
namespace {
-void buildTestVulkanRunnerPipeline(OpPassManager &passManager) {
+struct VulkanRunnerPipelineOptions
+ : public PassPipelineOptions<VulkanRunnerPipelineOptions> {
+ Option<bool> spirvWebGPUPrepare{
+ *this, "spirv-webgpu-prepare",
+ llvm::cl::desc("Run MLIR transforms used when targetting WebGPU")};
+};
+
+void buildTestVulkanRunnerPipeline(OpPassManager &passManager,
+ const VulkanRunnerPipelineOptions &options) {
passManager.addPass(createGpuKernelOutliningPass());
passManager.addPass(memref::createFoldMemRefAliasOpsPass());
+ GpuSPIRVAttachTargetOptions attachTargetOptions{};
+ attachTargetOptions.spirvVersion = "v1.0";
+ attachTargetOptions.spirvCapabilities.push_back("Shader");
+ attachTargetOptions.spirvExtensions.push_back(
+ "SPV_KHR_storage_buffer_storage_class");
+ passManager.addPass(createGpuSPIRVAttachTarget(attachTargetOptions));
+
ConvertToSPIRVPassOptions convertToSPIRVOptions{};
convertToSPIRVOptions.convertGPUModules = true;
+ convertToSPIRVOptions.nestInGPUModule = true;
passManager.addPass(createConvertToSPIRVPass(convertToSPIRVOptions));
- OpPassManager &modulePM = passManager.nest<spirv::ModuleOp>();
- modulePM.addPass(spirv::createSPIRVLowerABIAttributesPass());
- modulePM.addPass(spirv::createSPIRVUpdateVCEPass());
+
+ OpPassManager &gpuModulePM = passManager.nest<gpu::GPUModuleOp>();
+ OpPassManager &spirvModulePM = gpuModulePM.nest<spirv::ModuleOp>();
+ spirvModulePM.addPass(spirv::createSPIRVLowerABIAttributesPass());
+ spirvModulePM.addPass(spirv::createSPIRVUpdateVCEPass());
+ if (options.spirvWebGPUPrepare)
+ spirvModulePM.addPass(spirv::createSPIRVWebGPUPreparePass());
+
+ passManager.addPass(createGpuModuleToBinaryPass());
}
} // namespace
namespace mlir::test {
void registerTestVulkanRunnerPipeline() {
- PassPipelineRegistration<>(
+ PassPipelineRegistration<VulkanRunnerPipelineOptions>(
"test-vulkan-runner-pipeline",
"Runs a series of passes for lowering GPU-dialect MLIR to "
"SPIR-V-dialect MLIR intended for mlir-vulkan-runner.",
diff --git a/mlir/test/mlir-vulkan-runner/addui_extended.mlir b/mlir/test/mlir-vulkan-runner/addui_extended.mlir
index 158541f326be78..b8db4514214591 100644
--- a/mlir/test/mlir-vulkan-runner/addui_extended.mlir
+++ b/mlir/test/mlir-vulkan-runner/addui_extended.mlir
@@ -6,7 +6,7 @@
// RUN: --shared-libs=%vulkan-runtime-wrappers,%mlir_runner_utils \
// RUN: --entry-point-result=void | FileCheck %s
-// RUN: mlir-opt %s -test-vulkan-runner-pipeline -spirv-webgpu-prepare \
+// RUN: mlir-opt %s -test-vulkan-runner-pipeline=spirv-webgpu-prepare \
// RUN: | mlir-vulkan-runner - \
// RUN: --shared-libs=%vulkan-runtime-wrappers,%mlir_runner_utils \
// RUN: --entry-point-result=void | FileCheck %s
diff --git a/mlir/test/mlir-vulkan-runner/smul_extended.mlir b/mlir/test/mlir-vulkan-runner/smul_extended.mlir
index 2dd31d2ebb9a06..334aec843e1977 100644
--- a/mlir/test/mlir-vulkan-runner/smul_extended.mlir
+++ b/mlir/test/mlir-vulkan-runner/smul_extended.mlir
@@ -6,7 +6,7 @@
// RUN: --shared-libs=%vulkan-runtime-wrappers,%mlir_runner_utils \
// RUN: --entry-point-result=void | FileCheck %s
-// RUN: mlir-opt %s -test-vulkan-runner-pipeline -spirv-webgpu-prepare \
+// RUN: mlir-opt %s -test-vulkan-runner-pipeline=spirv-webgpu-prepare \
// RUN: | mlir-vulkan-runner - \
// RUN: --shared-libs=%vulkan-runtime-wrappers,%mlir_runner_utils \
// RUN: --entry-point-result=void | FileCheck %s
diff --git a/mlir/test/mlir-vulkan-runner/umul_extended.mlir b/mlir/test/mlir-vulkan-runner/umul_extended.mlir
index 78300d2fd81dd5..803b8c3d336d33 100644
--- a/mlir/test/mlir-vulkan-runner/umul_extended.mlir
+++ b/mlir/test/mlir-vulkan-runner/umul_extended.mlir
@@ -6,7 +6,7 @@
// RUN: --shared-libs=%vulkan-runtime-wrappers,%mlir_runner_utils \
// RUN: --entry-point-result=void | FileCheck %s
-// RUN: mlir-opt %s -test-vulkan-runner-pipeline -spirv-webgpu-prepare \
+// RUN: mlir-opt %s -test-vulkan-runner-pipeline=spirv-webgpu-prepare \
// RUN: | mlir-vulkan-runner - \
// RUN: --shared-libs=%vulkan-runtime-wrappers,%mlir_runner_utils \
// RUN: --entry-point-result=void | FileCheck %s
|
|
@llvm/pr-subscribers-mlir Author: Andrea Faulds (andfau-amd) ChangesThis commit is a further incremental step toward moving the whole mlir-vulkan-runner MLIR pass pipeline into mlir-opt (see #73457). The previous step was b225b3adf7b78387c9fcb97a3ff0e0a1e26eafe2, which moved all device passes prior to SPIR-V serialization into a new mlir-opt test pass, This commit changes how SPIR-V serialization is accomplished for Vulkan runner tests. Until now, this was done by the Vulkan-specific ConvertGpuLaunchFuncToVulkanLaunchFunc pass. With this commit, this responsibility is removed from that pass, and is instead done with the existing generic GpuModuleToBinaryPass. In addition, the SPIR-V serialization step is no longer done inside mlir-vulkan-runner, but rather inside mlir-opt (in the In order to enable this, this commit includes these supporting changes:
Full diff: https://github.com/llvm/llvm-project/pull/121494.diff 10 Files Affected:
diff --git a/mlir/include/mlir/Conversion/Passes.td b/mlir/include/mlir/Conversion/Passes.td
index 8835e0a9099fdd..e8713230814d01 100644
--- a/mlir/include/mlir/Conversion/Passes.td
+++ b/mlir/include/mlir/Conversion/Passes.td
@@ -61,7 +61,10 @@ def ConvertToSPIRVPass : Pass<"convert-to-spirv"> {
"Run vector unrolling to convert vector types in function bodies">,
Option<"convertGPUModules", "convert-gpu-modules", "bool",
/*default=*/"false",
- "Clone and convert GPU modules">
+ "Clone and convert GPU modules">,
+ Option<"nestInGPUModule", "nest-in-gpu-module", "bool",
+ /*default=*/"false",
+ "Put converted SPIR-V module inside the gpu.module instead of alongside it.">,
];
}
diff --git a/mlir/lib/Conversion/ConvertToSPIRV/ConvertToSPIRVPass.cpp b/mlir/lib/Conversion/ConvertToSPIRV/ConvertToSPIRVPass.cpp
index 4b7f7ff114deeb..ab9c048f561069 100644
--- a/mlir/lib/Conversion/ConvertToSPIRV/ConvertToSPIRVPass.cpp
+++ b/mlir/lib/Conversion/ConvertToSPIRV/ConvertToSPIRVPass.cpp
@@ -108,7 +108,10 @@ struct ConvertToSPIRVPass final
SmallVector<Operation *, 1> gpuModules;
OpBuilder builder(context);
op->walk([&](gpu::GPUModuleOp gpuModule) {
- builder.setInsertionPoint(gpuModule);
+ if (nestInGPUModule)
+ builder.setInsertionPointToStart(gpuModule.getBody());
+ else
+ builder.setInsertionPoint(gpuModule);
gpuModules.push_back(builder.clone(*gpuModule));
});
// Run conversion for each module independently as they can have
diff --git a/mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRVPass.cpp b/mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRVPass.cpp
index 08b451f7d5b325..509b6343057b99 100644
--- a/mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRVPass.cpp
+++ b/mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRVPass.cpp
@@ -71,7 +71,7 @@ void GPUToSPIRVPass::runOnOperation() {
// launch op still needs the original GPU kernel module.
// For Vulkan Shader capabilities, we insert the newly converted SPIR-V
// module right after the original GPU module, as that's the expectation of
- // the in-tree Vulkan runner.
+ // the in-tree SPIR-V CPU runner (the Vulkan runner does not use this pass).
// For OpenCL Kernel capabilities, we insert the newly converted SPIR-V
// module inside the original GPU module, as that's the expectaion of the
// normal GPU compilation pipeline.
diff --git a/mlir/lib/Conversion/GPUToVulkan/ConvertGPULaunchFuncToVulkanLaunchFunc.cpp b/mlir/lib/Conversion/GPUToVulkan/ConvertGPULaunchFuncToVulkanLaunchFunc.cpp
index 2d2251672230b6..69945cb6db8226 100644
--- a/mlir/lib/Conversion/GPUToVulkan/ConvertGPULaunchFuncToVulkanLaunchFunc.cpp
+++ b/mlir/lib/Conversion/GPUToVulkan/ConvertGPULaunchFuncToVulkanLaunchFunc.cpp
@@ -7,9 +7,8 @@
//===----------------------------------------------------------------------===//
//
// This file implements a pass to convert gpu launch function into a vulkan
-// launch function. Creates a SPIR-V binary shader from the `spirv::ModuleOp`
-// using `spirv::serialize` function, attaches binary data and entry point name
-// as an attributes to vulkan launch call op.
+// launch function. Extracts the SPIR-V from a `gpu::BinaryOp` and attaches it
+// along with the entry point name as attributes to a Vulkan launch call op.
//
//===----------------------------------------------------------------------===//
@@ -40,10 +39,9 @@ static constexpr const char *kVulkanLaunch = "vulkanLaunch";
namespace {
-/// A pass to convert gpu launch op to vulkan launch call op, by creating a
-/// SPIR-V binary shader from `spirv::ModuleOp` using `spirv::serialize`
-/// function and attaching binary data and entry point name as an attributes to
-/// created vulkan launch call op.
+/// A pass to convert gpu launch op to vulkan launch call op, by extracting a
+/// SPIR-V binary shader from a `gpu::BinaryOp` and attaching binary data and
+/// entry point name as an attributes to created vulkan launch call op.
class ConvertGpuLaunchFuncToVulkanLaunchFunc
: public impl::ConvertGpuLaunchFuncToVulkanLaunchFuncBase<
ConvertGpuLaunchFuncToVulkanLaunchFunc> {
@@ -51,10 +49,10 @@ class ConvertGpuLaunchFuncToVulkanLaunchFunc
void runOnOperation() override;
private:
- /// Creates a SPIR-V binary shader from the given `module` using
- /// `spirv::serialize` function.
- LogicalResult createBinaryShader(ModuleOp module,
- std::vector<char> &binaryShader);
+ /// Extracts a SPIR-V binary shader from the given `module`, if any.
+ /// Note that this also removes the binary from the IR.
+ LogicalResult getBinaryShader(ModuleOp module,
+ std::vector<char> &binaryShader);
/// Converts the given `launchOp` to vulkan launch call.
void convertGpuLaunchFunc(gpu::LaunchFuncOp launchOp);
@@ -135,21 +133,38 @@ LogicalResult ConvertGpuLaunchFuncToVulkanLaunchFunc::declareVulkanLaunchFunc(
return success();
}
-LogicalResult ConvertGpuLaunchFuncToVulkanLaunchFunc::createBinaryShader(
+LogicalResult ConvertGpuLaunchFuncToVulkanLaunchFunc::getBinaryShader(
ModuleOp module, std::vector<char> &binaryShader) {
bool done = false;
SmallVector<uint32_t, 0> binary;
- for (auto spirvModule : module.getOps<spirv::ModuleOp>()) {
+ gpu::BinaryOp *binaryToErase;
+ for (auto gpuBinary : module.getOps<gpu::BinaryOp>()) {
if (done)
- return spirvModule.emitError("should only contain one 'spirv.module' op");
+ return gpuBinary.emitError("should only contain one 'gpu.binary' op");
done = true;
- if (failed(spirv::serialize(spirvModule, binary)))
- return failure();
+ ArrayRef<Attribute> objects = gpuBinary.getObjectsAttr().getValue();
+ if (objects.size() != 1)
+ return gpuBinary.emitError("should only contain a single object");
+
+ auto object = cast<gpu::ObjectAttr>(objects[0]);
+
+ if (!isa<spirv::TargetEnvAttr>(object.getTarget()))
+ return gpuBinary.emitError(
+ "should contain an object with a SPIR-V target environment");
+
+ StringAttr objectStrAttr = object.getObject();
+ StringRef objectStr = objectStrAttr.getValue();
+ binaryShader.insert(binaryShader.end(), objectStr.bytes_begin(),
+ objectStr.bytes_end());
+
+ binaryToErase = &gpuBinary;
}
- binaryShader.resize(binary.size() * sizeof(uint32_t));
- std::memcpy(binaryShader.data(), reinterpret_cast<char *>(binary.data()),
- binaryShader.size());
+ if (!done)
+ return module.emitError("should contain a 'gpu.binary' op");
+
+ // Remove the binary to avoid confusing later conversion passes.
+ binaryToErase->erase();
return success();
}
@@ -159,9 +174,9 @@ void ConvertGpuLaunchFuncToVulkanLaunchFunc::convertGpuLaunchFunc(
OpBuilder builder(launchOp);
Location loc = launchOp.getLoc();
- // Serialize `spirv::Module` into binary form.
std::vector<char> binary;
- if (failed(createBinaryShader(module, binary)))
+ // Extract SPIR-V from `gpu.binary` op.
+ if (failed(getBinaryShader(module, binary)))
return signalPassFailure();
// Declare vulkan launch function.
diff --git a/mlir/test/Conversion/ConvertToSPIRV/convert-gpu-modules-nested.mlir b/mlir/test/Conversion/ConvertToSPIRV/convert-gpu-modules-nested.mlir
new file mode 100644
index 00000000000000..33fa0f859a5c77
--- /dev/null
+++ b/mlir/test/Conversion/ConvertToSPIRV/convert-gpu-modules-nested.mlir
@@ -0,0 +1,30 @@
+// RUN: mlir-opt -convert-to-spirv="convert-gpu-modules=true nest-in-gpu-module=true run-signature-conversion=false run-vector-unrolling=false" -split-input-file %s | FileCheck %s
+
+module attributes {
+ gpu.container_module,
+ spirv.target_env = #spirv.target_env<#spirv.vce<v1.0, [Shader], []>, #spirv.resource_limits<>>
+} {
+ // CHECK-LABEL: func.func @main
+ // CHECK: %[[C1:.*]] = arith.constant 1 : index
+ // CHECK: gpu.launch_func @[[$KERNELS_1:.*]]::@[[$BUILTIN_WG_ID_X:.*]] blocks in (%[[C1]], %[[C1]], %[[C1]]) threads in (%[[C1]], %[[C1]], %[[C1]])
+ func.func @main() {
+ %c1 = arith.constant 1 : index
+ gpu.launch_func @kernels_1::@builtin_workgroup_id_x
+ blocks in (%c1, %c1, %c1) threads in (%c1, %c1, %c1)
+ return
+ }
+
+ // CHECK: gpu.module @[[$KERNELS_1]]
+ // CHECK: spirv.module @{{.*}} Logical GLSL450
+ // CHECK: spirv.func @[[$BUILTIN_WG_ID_X]]
+ // CHECK: spirv.mlir.addressof
+ // CHECK: spirv.Load "Input"
+ // CHECK: spirv.CompositeExtract
+ gpu.module @kernels_1 {
+ gpu.func @builtin_workgroup_id_x() kernel
+ attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {
+ %0 = gpu.block_id x
+ gpu.return
+ }
+ }
+}
diff --git a/mlir/test/Conversion/GPUToVulkan/lower-gpu-launch-vulkan-launch.mlir b/mlir/test/Conversion/GPUToVulkan/lower-gpu-launch-vulkan-launch.mlir
index 665d0a33abedc2..96ee1866517e6d 100644
--- a/mlir/test/Conversion/GPUToVulkan/lower-gpu-launch-vulkan-launch.mlir
+++ b/mlir/test/Conversion/GPUToVulkan/lower-gpu-launch-vulkan-launch.mlir
@@ -1,24 +1,24 @@
-// RUN: mlir-opt %s -convert-gpu-launch-to-vulkan-launch | FileCheck %s
+// RUN: mlir-opt %s -pass-pipeline='builtin.module(spirv-attach-target{ver=v1.0 caps=Shader exts=SPV_KHR_storage_buffer_storage_class},gpu-module-to-binary,convert-gpu-launch-to-vulkan-launch)' | FileCheck %s
// CHECK: %[[resource:.*]] = memref.alloc() : memref<12xf32>
// CHECK: %[[index:.*]] = arith.constant 1 : index
// CHECK: call @vulkanLaunch(%[[index]], %[[index]], %[[index]], %[[resource]]) {spirv_blob = "{{.*}}", spirv_element_types = [f32], spirv_entry_point = "kernel"}
module attributes {gpu.container_module} {
- spirv.module Logical GLSL450 requires #spirv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]> {
- spirv.GlobalVariable @kernel_arg_0 bind(0, 0) : !spirv.ptr<!spirv.struct<(!spirv.array<12 x f32, stride=4> [0])>, StorageBuffer>
- spirv.func @kernel() "None" attributes {workgroup_attributions = 0 : i64} {
- %0 = spirv.mlir.addressof @kernel_arg_0 : !spirv.ptr<!spirv.struct<(!spirv.array<12 x f32, stride=4> [0])>, StorageBuffer>
- %2 = spirv.Constant 0 : i32
- %3 = spirv.mlir.addressof @kernel_arg_0 : !spirv.ptr<!spirv.struct<(!spirv.array<12 x f32, stride=4> [0])>, StorageBuffer>
- %4 = spirv.AccessChain %0[%2, %2] : !spirv.ptr<!spirv.struct<(!spirv.array<12 x f32, stride=4> [0])>, StorageBuffer>, i32, i32 -> !spirv.ptr<f32, StorageBuffer>
- %5 = spirv.Load "StorageBuffer" %4 : f32
- spirv.Return
- }
- spirv.EntryPoint "GLCompute" @kernel
- spirv.ExecutionMode @kernel "LocalSize", 1, 1, 1
- }
gpu.module @kernels {
+ spirv.module Logical GLSL450 requires #spirv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]> {
+ spirv.GlobalVariable @kernel_arg_0 bind(0, 0) : !spirv.ptr<!spirv.struct<(!spirv.array<12 x f32, stride=4> [0])>, StorageBuffer>
+ spirv.func @kernel() "None" attributes {workgroup_attributions = 0 : i64} {
+ %0 = spirv.mlir.addressof @kernel_arg_0 : !spirv.ptr<!spirv.struct<(!spirv.array<12 x f32, stride=4> [0])>, StorageBuffer>
+ %2 = spirv.Constant 0 : i32
+ %3 = spirv.mlir.addressof @kernel_arg_0 : !spirv.ptr<!spirv.struct<(!spirv.array<12 x f32, stride=4> [0])>, StorageBuffer>
+ %4 = spirv.AccessChain %0[%2, %2] : !spirv.ptr<!spirv.struct<(!spirv.array<12 x f32, stride=4> [0])>, StorageBuffer>, i32, i32 -> !spirv.ptr<f32, StorageBuffer>
+ %5 = spirv.Load "StorageBuffer" %4 : f32
+ spirv.Return
+ }
+ spirv.EntryPoint "GLCompute" @kernel
+ spirv.ExecutionMode @kernel "LocalSize", 1, 1, 1
+ }
gpu.func @kernel(%arg0: memref<12xf32>) kernel {
gpu.return
}
diff --git a/mlir/test/lib/Pass/TestVulkanRunnerPipeline.cpp b/mlir/test/lib/Pass/TestVulkanRunnerPipeline.cpp
index eda9aa9f9efef7..9bd4c42a1cdfb5 100644
--- a/mlir/test/lib/Pass/TestVulkanRunnerPipeline.cpp
+++ b/mlir/test/lib/Pass/TestVulkanRunnerPipeline.cpp
@@ -12,33 +12,57 @@
#include "mlir/Conversion/ConvertToSPIRV/ConvertToSPIRVPass.h"
#include "mlir/Conversion/GPUToSPIRV/GPUToSPIRVPass.h"
+#include "mlir/Dialect/GPU/IR/GPUDialect.h"
#include "mlir/Dialect/GPU/Transforms/Passes.h"
#include "mlir/Dialect/MemRef/Transforms/Passes.h"
#include "mlir/Dialect/SPIRV/IR/SPIRVOps.h"
#include "mlir/Dialect/SPIRV/Transforms/Passes.h"
#include "mlir/Pass/PassManager.h"
+#include "mlir/Pass/PassOptions.h"
using namespace mlir;
namespace {
-void buildTestVulkanRunnerPipeline(OpPassManager &passManager) {
+struct VulkanRunnerPipelineOptions
+ : public PassPipelineOptions<VulkanRunnerPipelineOptions> {
+ Option<bool> spirvWebGPUPrepare{
+ *this, "spirv-webgpu-prepare",
+ llvm::cl::desc("Run MLIR transforms used when targetting WebGPU")};
+};
+
+void buildTestVulkanRunnerPipeline(OpPassManager &passManager,
+ const VulkanRunnerPipelineOptions &options) {
passManager.addPass(createGpuKernelOutliningPass());
passManager.addPass(memref::createFoldMemRefAliasOpsPass());
+ GpuSPIRVAttachTargetOptions attachTargetOptions{};
+ attachTargetOptions.spirvVersion = "v1.0";
+ attachTargetOptions.spirvCapabilities.push_back("Shader");
+ attachTargetOptions.spirvExtensions.push_back(
+ "SPV_KHR_storage_buffer_storage_class");
+ passManager.addPass(createGpuSPIRVAttachTarget(attachTargetOptions));
+
ConvertToSPIRVPassOptions convertToSPIRVOptions{};
convertToSPIRVOptions.convertGPUModules = true;
+ convertToSPIRVOptions.nestInGPUModule = true;
passManager.addPass(createConvertToSPIRVPass(convertToSPIRVOptions));
- OpPassManager &modulePM = passManager.nest<spirv::ModuleOp>();
- modulePM.addPass(spirv::createSPIRVLowerABIAttributesPass());
- modulePM.addPass(spirv::createSPIRVUpdateVCEPass());
+
+ OpPassManager &gpuModulePM = passManager.nest<gpu::GPUModuleOp>();
+ OpPassManager &spirvModulePM = gpuModulePM.nest<spirv::ModuleOp>();
+ spirvModulePM.addPass(spirv::createSPIRVLowerABIAttributesPass());
+ spirvModulePM.addPass(spirv::createSPIRVUpdateVCEPass());
+ if (options.spirvWebGPUPrepare)
+ spirvModulePM.addPass(spirv::createSPIRVWebGPUPreparePass());
+
+ passManager.addPass(createGpuModuleToBinaryPass());
}
} // namespace
namespace mlir::test {
void registerTestVulkanRunnerPipeline() {
- PassPipelineRegistration<>(
+ PassPipelineRegistration<VulkanRunnerPipelineOptions>(
"test-vulkan-runner-pipeline",
"Runs a series of passes for lowering GPU-dialect MLIR to "
"SPIR-V-dialect MLIR intended for mlir-vulkan-runner.",
diff --git a/mlir/test/mlir-vulkan-runner/addui_extended.mlir b/mlir/test/mlir-vulkan-runner/addui_extended.mlir
index 158541f326be78..b8db4514214591 100644
--- a/mlir/test/mlir-vulkan-runner/addui_extended.mlir
+++ b/mlir/test/mlir-vulkan-runner/addui_extended.mlir
@@ -6,7 +6,7 @@
// RUN: --shared-libs=%vulkan-runtime-wrappers,%mlir_runner_utils \
// RUN: --entry-point-result=void | FileCheck %s
-// RUN: mlir-opt %s -test-vulkan-runner-pipeline -spirv-webgpu-prepare \
+// RUN: mlir-opt %s -test-vulkan-runner-pipeline=spirv-webgpu-prepare \
// RUN: | mlir-vulkan-runner - \
// RUN: --shared-libs=%vulkan-runtime-wrappers,%mlir_runner_utils \
// RUN: --entry-point-result=void | FileCheck %s
diff --git a/mlir/test/mlir-vulkan-runner/smul_extended.mlir b/mlir/test/mlir-vulkan-runner/smul_extended.mlir
index 2dd31d2ebb9a06..334aec843e1977 100644
--- a/mlir/test/mlir-vulkan-runner/smul_extended.mlir
+++ b/mlir/test/mlir-vulkan-runner/smul_extended.mlir
@@ -6,7 +6,7 @@
// RUN: --shared-libs=%vulkan-runtime-wrappers,%mlir_runner_utils \
// RUN: --entry-point-result=void | FileCheck %s
-// RUN: mlir-opt %s -test-vulkan-runner-pipeline -spirv-webgpu-prepare \
+// RUN: mlir-opt %s -test-vulkan-runner-pipeline=spirv-webgpu-prepare \
// RUN: | mlir-vulkan-runner - \
// RUN: --shared-libs=%vulkan-runtime-wrappers,%mlir_runner_utils \
// RUN: --entry-point-result=void | FileCheck %s
diff --git a/mlir/test/mlir-vulkan-runner/umul_extended.mlir b/mlir/test/mlir-vulkan-runner/umul_extended.mlir
index 78300d2fd81dd5..803b8c3d336d33 100644
--- a/mlir/test/mlir-vulkan-runner/umul_extended.mlir
+++ b/mlir/test/mlir-vulkan-runner/umul_extended.mlir
@@ -6,7 +6,7 @@
// RUN: --shared-libs=%vulkan-runtime-wrappers,%mlir_runner_utils \
// RUN: --entry-point-result=void | FileCheck %s
-// RUN: mlir-opt %s -test-vulkan-runner-pipeline -spirv-webgpu-prepare \
+// RUN: mlir-opt %s -test-vulkan-runner-pipeline=spirv-webgpu-prepare \
// RUN: | mlir-vulkan-runner - \
// RUN: --shared-libs=%vulkan-runtime-wrappers,%mlir_runner_utils \
// RUN: --entry-point-result=void | FileCheck %s
|
|
@llvm/pr-subscribers-mlir-spirv Author: Andrea Faulds (andfau-amd) ChangesThis commit is a further incremental step toward moving the whole mlir-vulkan-runner MLIR pass pipeline into mlir-opt (see #73457). The previous step was b225b3adf7b78387c9fcb97a3ff0e0a1e26eafe2, which moved all device passes prior to SPIR-V serialization into a new mlir-opt test pass, This commit changes how SPIR-V serialization is accomplished for Vulkan runner tests. Until now, this was done by the Vulkan-specific ConvertGpuLaunchFuncToVulkanLaunchFunc pass. With this commit, this responsibility is removed from that pass, and is instead done with the existing generic GpuModuleToBinaryPass. In addition, the SPIR-V serialization step is no longer done inside mlir-vulkan-runner, but rather inside mlir-opt (in the In order to enable this, this commit includes these supporting changes:
Full diff: https://github.com/llvm/llvm-project/pull/121494.diff 10 Files Affected:
diff --git a/mlir/include/mlir/Conversion/Passes.td b/mlir/include/mlir/Conversion/Passes.td
index 8835e0a9099fdd..e8713230814d01 100644
--- a/mlir/include/mlir/Conversion/Passes.td
+++ b/mlir/include/mlir/Conversion/Passes.td
@@ -61,7 +61,10 @@ def ConvertToSPIRVPass : Pass<"convert-to-spirv"> {
"Run vector unrolling to convert vector types in function bodies">,
Option<"convertGPUModules", "convert-gpu-modules", "bool",
/*default=*/"false",
- "Clone and convert GPU modules">
+ "Clone and convert GPU modules">,
+ Option<"nestInGPUModule", "nest-in-gpu-module", "bool",
+ /*default=*/"false",
+ "Put converted SPIR-V module inside the gpu.module instead of alongside it.">,
];
}
diff --git a/mlir/lib/Conversion/ConvertToSPIRV/ConvertToSPIRVPass.cpp b/mlir/lib/Conversion/ConvertToSPIRV/ConvertToSPIRVPass.cpp
index 4b7f7ff114deeb..ab9c048f561069 100644
--- a/mlir/lib/Conversion/ConvertToSPIRV/ConvertToSPIRVPass.cpp
+++ b/mlir/lib/Conversion/ConvertToSPIRV/ConvertToSPIRVPass.cpp
@@ -108,7 +108,10 @@ struct ConvertToSPIRVPass final
SmallVector<Operation *, 1> gpuModules;
OpBuilder builder(context);
op->walk([&](gpu::GPUModuleOp gpuModule) {
- builder.setInsertionPoint(gpuModule);
+ if (nestInGPUModule)
+ builder.setInsertionPointToStart(gpuModule.getBody());
+ else
+ builder.setInsertionPoint(gpuModule);
gpuModules.push_back(builder.clone(*gpuModule));
});
// Run conversion for each module independently as they can have
diff --git a/mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRVPass.cpp b/mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRVPass.cpp
index 08b451f7d5b325..509b6343057b99 100644
--- a/mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRVPass.cpp
+++ b/mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRVPass.cpp
@@ -71,7 +71,7 @@ void GPUToSPIRVPass::runOnOperation() {
// launch op still needs the original GPU kernel module.
// For Vulkan Shader capabilities, we insert the newly converted SPIR-V
// module right after the original GPU module, as that's the expectation of
- // the in-tree Vulkan runner.
+ // the in-tree SPIR-V CPU runner (the Vulkan runner does not use this pass).
// For OpenCL Kernel capabilities, we insert the newly converted SPIR-V
// module inside the original GPU module, as that's the expectaion of the
// normal GPU compilation pipeline.
diff --git a/mlir/lib/Conversion/GPUToVulkan/ConvertGPULaunchFuncToVulkanLaunchFunc.cpp b/mlir/lib/Conversion/GPUToVulkan/ConvertGPULaunchFuncToVulkanLaunchFunc.cpp
index 2d2251672230b6..69945cb6db8226 100644
--- a/mlir/lib/Conversion/GPUToVulkan/ConvertGPULaunchFuncToVulkanLaunchFunc.cpp
+++ b/mlir/lib/Conversion/GPUToVulkan/ConvertGPULaunchFuncToVulkanLaunchFunc.cpp
@@ -7,9 +7,8 @@
//===----------------------------------------------------------------------===//
//
// This file implements a pass to convert gpu launch function into a vulkan
-// launch function. Creates a SPIR-V binary shader from the `spirv::ModuleOp`
-// using `spirv::serialize` function, attaches binary data and entry point name
-// as an attributes to vulkan launch call op.
+// launch function. Extracts the SPIR-V from a `gpu::BinaryOp` and attaches it
+// along with the entry point name as attributes to a Vulkan launch call op.
//
//===----------------------------------------------------------------------===//
@@ -40,10 +39,9 @@ static constexpr const char *kVulkanLaunch = "vulkanLaunch";
namespace {
-/// A pass to convert gpu launch op to vulkan launch call op, by creating a
-/// SPIR-V binary shader from `spirv::ModuleOp` using `spirv::serialize`
-/// function and attaching binary data and entry point name as an attributes to
-/// created vulkan launch call op.
+/// A pass to convert gpu launch op to vulkan launch call op, by extracting a
+/// SPIR-V binary shader from a `gpu::BinaryOp` and attaching binary data and
+/// entry point name as an attributes to created vulkan launch call op.
class ConvertGpuLaunchFuncToVulkanLaunchFunc
: public impl::ConvertGpuLaunchFuncToVulkanLaunchFuncBase<
ConvertGpuLaunchFuncToVulkanLaunchFunc> {
@@ -51,10 +49,10 @@ class ConvertGpuLaunchFuncToVulkanLaunchFunc
void runOnOperation() override;
private:
- /// Creates a SPIR-V binary shader from the given `module` using
- /// `spirv::serialize` function.
- LogicalResult createBinaryShader(ModuleOp module,
- std::vector<char> &binaryShader);
+ /// Extracts a SPIR-V binary shader from the given `module`, if any.
+ /// Note that this also removes the binary from the IR.
+ LogicalResult getBinaryShader(ModuleOp module,
+ std::vector<char> &binaryShader);
/// Converts the given `launchOp` to vulkan launch call.
void convertGpuLaunchFunc(gpu::LaunchFuncOp launchOp);
@@ -135,21 +133,38 @@ LogicalResult ConvertGpuLaunchFuncToVulkanLaunchFunc::declareVulkanLaunchFunc(
return success();
}
-LogicalResult ConvertGpuLaunchFuncToVulkanLaunchFunc::createBinaryShader(
+LogicalResult ConvertGpuLaunchFuncToVulkanLaunchFunc::getBinaryShader(
ModuleOp module, std::vector<char> &binaryShader) {
bool done = false;
SmallVector<uint32_t, 0> binary;
- for (auto spirvModule : module.getOps<spirv::ModuleOp>()) {
+ gpu::BinaryOp *binaryToErase;
+ for (auto gpuBinary : module.getOps<gpu::BinaryOp>()) {
if (done)
- return spirvModule.emitError("should only contain one 'spirv.module' op");
+ return gpuBinary.emitError("should only contain one 'gpu.binary' op");
done = true;
- if (failed(spirv::serialize(spirvModule, binary)))
- return failure();
+ ArrayRef<Attribute> objects = gpuBinary.getObjectsAttr().getValue();
+ if (objects.size() != 1)
+ return gpuBinary.emitError("should only contain a single object");
+
+ auto object = cast<gpu::ObjectAttr>(objects[0]);
+
+ if (!isa<spirv::TargetEnvAttr>(object.getTarget()))
+ return gpuBinary.emitError(
+ "should contain an object with a SPIR-V target environment");
+
+ StringAttr objectStrAttr = object.getObject();
+ StringRef objectStr = objectStrAttr.getValue();
+ binaryShader.insert(binaryShader.end(), objectStr.bytes_begin(),
+ objectStr.bytes_end());
+
+ binaryToErase = &gpuBinary;
}
- binaryShader.resize(binary.size() * sizeof(uint32_t));
- std::memcpy(binaryShader.data(), reinterpret_cast<char *>(binary.data()),
- binaryShader.size());
+ if (!done)
+ return module.emitError("should contain a 'gpu.binary' op");
+
+ // Remove the binary to avoid confusing later conversion passes.
+ binaryToErase->erase();
return success();
}
@@ -159,9 +174,9 @@ void ConvertGpuLaunchFuncToVulkanLaunchFunc::convertGpuLaunchFunc(
OpBuilder builder(launchOp);
Location loc = launchOp.getLoc();
- // Serialize `spirv::Module` into binary form.
std::vector<char> binary;
- if (failed(createBinaryShader(module, binary)))
+ // Extract SPIR-V from `gpu.binary` op.
+ if (failed(getBinaryShader(module, binary)))
return signalPassFailure();
// Declare vulkan launch function.
diff --git a/mlir/test/Conversion/ConvertToSPIRV/convert-gpu-modules-nested.mlir b/mlir/test/Conversion/ConvertToSPIRV/convert-gpu-modules-nested.mlir
new file mode 100644
index 00000000000000..33fa0f859a5c77
--- /dev/null
+++ b/mlir/test/Conversion/ConvertToSPIRV/convert-gpu-modules-nested.mlir
@@ -0,0 +1,30 @@
+// RUN: mlir-opt -convert-to-spirv="convert-gpu-modules=true nest-in-gpu-module=true run-signature-conversion=false run-vector-unrolling=false" -split-input-file %s | FileCheck %s
+
+module attributes {
+ gpu.container_module,
+ spirv.target_env = #spirv.target_env<#spirv.vce<v1.0, [Shader], []>, #spirv.resource_limits<>>
+} {
+ // CHECK-LABEL: func.func @main
+ // CHECK: %[[C1:.*]] = arith.constant 1 : index
+ // CHECK: gpu.launch_func @[[$KERNELS_1:.*]]::@[[$BUILTIN_WG_ID_X:.*]] blocks in (%[[C1]], %[[C1]], %[[C1]]) threads in (%[[C1]], %[[C1]], %[[C1]])
+ func.func @main() {
+ %c1 = arith.constant 1 : index
+ gpu.launch_func @kernels_1::@builtin_workgroup_id_x
+ blocks in (%c1, %c1, %c1) threads in (%c1, %c1, %c1)
+ return
+ }
+
+ // CHECK: gpu.module @[[$KERNELS_1]]
+ // CHECK: spirv.module @{{.*}} Logical GLSL450
+ // CHECK: spirv.func @[[$BUILTIN_WG_ID_X]]
+ // CHECK: spirv.mlir.addressof
+ // CHECK: spirv.Load "Input"
+ // CHECK: spirv.CompositeExtract
+ gpu.module @kernels_1 {
+ gpu.func @builtin_workgroup_id_x() kernel
+ attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {
+ %0 = gpu.block_id x
+ gpu.return
+ }
+ }
+}
diff --git a/mlir/test/Conversion/GPUToVulkan/lower-gpu-launch-vulkan-launch.mlir b/mlir/test/Conversion/GPUToVulkan/lower-gpu-launch-vulkan-launch.mlir
index 665d0a33abedc2..96ee1866517e6d 100644
--- a/mlir/test/Conversion/GPUToVulkan/lower-gpu-launch-vulkan-launch.mlir
+++ b/mlir/test/Conversion/GPUToVulkan/lower-gpu-launch-vulkan-launch.mlir
@@ -1,24 +1,24 @@
-// RUN: mlir-opt %s -convert-gpu-launch-to-vulkan-launch | FileCheck %s
+// RUN: mlir-opt %s -pass-pipeline='builtin.module(spirv-attach-target{ver=v1.0 caps=Shader exts=SPV_KHR_storage_buffer_storage_class},gpu-module-to-binary,convert-gpu-launch-to-vulkan-launch)' | FileCheck %s
// CHECK: %[[resource:.*]] = memref.alloc() : memref<12xf32>
// CHECK: %[[index:.*]] = arith.constant 1 : index
// CHECK: call @vulkanLaunch(%[[index]], %[[index]], %[[index]], %[[resource]]) {spirv_blob = "{{.*}}", spirv_element_types = [f32], spirv_entry_point = "kernel"}
module attributes {gpu.container_module} {
- spirv.module Logical GLSL450 requires #spirv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]> {
- spirv.GlobalVariable @kernel_arg_0 bind(0, 0) : !spirv.ptr<!spirv.struct<(!spirv.array<12 x f32, stride=4> [0])>, StorageBuffer>
- spirv.func @kernel() "None" attributes {workgroup_attributions = 0 : i64} {
- %0 = spirv.mlir.addressof @kernel_arg_0 : !spirv.ptr<!spirv.struct<(!spirv.array<12 x f32, stride=4> [0])>, StorageBuffer>
- %2 = spirv.Constant 0 : i32
- %3 = spirv.mlir.addressof @kernel_arg_0 : !spirv.ptr<!spirv.struct<(!spirv.array<12 x f32, stride=4> [0])>, StorageBuffer>
- %4 = spirv.AccessChain %0[%2, %2] : !spirv.ptr<!spirv.struct<(!spirv.array<12 x f32, stride=4> [0])>, StorageBuffer>, i32, i32 -> !spirv.ptr<f32, StorageBuffer>
- %5 = spirv.Load "StorageBuffer" %4 : f32
- spirv.Return
- }
- spirv.EntryPoint "GLCompute" @kernel
- spirv.ExecutionMode @kernel "LocalSize", 1, 1, 1
- }
gpu.module @kernels {
+ spirv.module Logical GLSL450 requires #spirv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]> {
+ spirv.GlobalVariable @kernel_arg_0 bind(0, 0) : !spirv.ptr<!spirv.struct<(!spirv.array<12 x f32, stride=4> [0])>, StorageBuffer>
+ spirv.func @kernel() "None" attributes {workgroup_attributions = 0 : i64} {
+ %0 = spirv.mlir.addressof @kernel_arg_0 : !spirv.ptr<!spirv.struct<(!spirv.array<12 x f32, stride=4> [0])>, StorageBuffer>
+ %2 = spirv.Constant 0 : i32
+ %3 = spirv.mlir.addressof @kernel_arg_0 : !spirv.ptr<!spirv.struct<(!spirv.array<12 x f32, stride=4> [0])>, StorageBuffer>
+ %4 = spirv.AccessChain %0[%2, %2] : !spirv.ptr<!spirv.struct<(!spirv.array<12 x f32, stride=4> [0])>, StorageBuffer>, i32, i32 -> !spirv.ptr<f32, StorageBuffer>
+ %5 = spirv.Load "StorageBuffer" %4 : f32
+ spirv.Return
+ }
+ spirv.EntryPoint "GLCompute" @kernel
+ spirv.ExecutionMode @kernel "LocalSize", 1, 1, 1
+ }
gpu.func @kernel(%arg0: memref<12xf32>) kernel {
gpu.return
}
diff --git a/mlir/test/lib/Pass/TestVulkanRunnerPipeline.cpp b/mlir/test/lib/Pass/TestVulkanRunnerPipeline.cpp
index eda9aa9f9efef7..9bd4c42a1cdfb5 100644
--- a/mlir/test/lib/Pass/TestVulkanRunnerPipeline.cpp
+++ b/mlir/test/lib/Pass/TestVulkanRunnerPipeline.cpp
@@ -12,33 +12,57 @@
#include "mlir/Conversion/ConvertToSPIRV/ConvertToSPIRVPass.h"
#include "mlir/Conversion/GPUToSPIRV/GPUToSPIRVPass.h"
+#include "mlir/Dialect/GPU/IR/GPUDialect.h"
#include "mlir/Dialect/GPU/Transforms/Passes.h"
#include "mlir/Dialect/MemRef/Transforms/Passes.h"
#include "mlir/Dialect/SPIRV/IR/SPIRVOps.h"
#include "mlir/Dialect/SPIRV/Transforms/Passes.h"
#include "mlir/Pass/PassManager.h"
+#include "mlir/Pass/PassOptions.h"
using namespace mlir;
namespace {
-void buildTestVulkanRunnerPipeline(OpPassManager &passManager) {
+struct VulkanRunnerPipelineOptions
+ : public PassPipelineOptions<VulkanRunnerPipelineOptions> {
+ Option<bool> spirvWebGPUPrepare{
+ *this, "spirv-webgpu-prepare",
+ llvm::cl::desc("Run MLIR transforms used when targetting WebGPU")};
+};
+
+void buildTestVulkanRunnerPipeline(OpPassManager &passManager,
+ const VulkanRunnerPipelineOptions &options) {
passManager.addPass(createGpuKernelOutliningPass());
passManager.addPass(memref::createFoldMemRefAliasOpsPass());
+ GpuSPIRVAttachTargetOptions attachTargetOptions{};
+ attachTargetOptions.spirvVersion = "v1.0";
+ attachTargetOptions.spirvCapabilities.push_back("Shader");
+ attachTargetOptions.spirvExtensions.push_back(
+ "SPV_KHR_storage_buffer_storage_class");
+ passManager.addPass(createGpuSPIRVAttachTarget(attachTargetOptions));
+
ConvertToSPIRVPassOptions convertToSPIRVOptions{};
convertToSPIRVOptions.convertGPUModules = true;
+ convertToSPIRVOptions.nestInGPUModule = true;
passManager.addPass(createConvertToSPIRVPass(convertToSPIRVOptions));
- OpPassManager &modulePM = passManager.nest<spirv::ModuleOp>();
- modulePM.addPass(spirv::createSPIRVLowerABIAttributesPass());
- modulePM.addPass(spirv::createSPIRVUpdateVCEPass());
+
+ OpPassManager &gpuModulePM = passManager.nest<gpu::GPUModuleOp>();
+ OpPassManager &spirvModulePM = gpuModulePM.nest<spirv::ModuleOp>();
+ spirvModulePM.addPass(spirv::createSPIRVLowerABIAttributesPass());
+ spirvModulePM.addPass(spirv::createSPIRVUpdateVCEPass());
+ if (options.spirvWebGPUPrepare)
+ spirvModulePM.addPass(spirv::createSPIRVWebGPUPreparePass());
+
+ passManager.addPass(createGpuModuleToBinaryPass());
}
} // namespace
namespace mlir::test {
void registerTestVulkanRunnerPipeline() {
- PassPipelineRegistration<>(
+ PassPipelineRegistration<VulkanRunnerPipelineOptions>(
"test-vulkan-runner-pipeline",
"Runs a series of passes for lowering GPU-dialect MLIR to "
"SPIR-V-dialect MLIR intended for mlir-vulkan-runner.",
diff --git a/mlir/test/mlir-vulkan-runner/addui_extended.mlir b/mlir/test/mlir-vulkan-runner/addui_extended.mlir
index 158541f326be78..b8db4514214591 100644
--- a/mlir/test/mlir-vulkan-runner/addui_extended.mlir
+++ b/mlir/test/mlir-vulkan-runner/addui_extended.mlir
@@ -6,7 +6,7 @@
// RUN: --shared-libs=%vulkan-runtime-wrappers,%mlir_runner_utils \
// RUN: --entry-point-result=void | FileCheck %s
-// RUN: mlir-opt %s -test-vulkan-runner-pipeline -spirv-webgpu-prepare \
+// RUN: mlir-opt %s -test-vulkan-runner-pipeline=spirv-webgpu-prepare \
// RUN: | mlir-vulkan-runner - \
// RUN: --shared-libs=%vulkan-runtime-wrappers,%mlir_runner_utils \
// RUN: --entry-point-result=void | FileCheck %s
diff --git a/mlir/test/mlir-vulkan-runner/smul_extended.mlir b/mlir/test/mlir-vulkan-runner/smul_extended.mlir
index 2dd31d2ebb9a06..334aec843e1977 100644
--- a/mlir/test/mlir-vulkan-runner/smul_extended.mlir
+++ b/mlir/test/mlir-vulkan-runner/smul_extended.mlir
@@ -6,7 +6,7 @@
// RUN: --shared-libs=%vulkan-runtime-wrappers,%mlir_runner_utils \
// RUN: --entry-point-result=void | FileCheck %s
-// RUN: mlir-opt %s -test-vulkan-runner-pipeline -spirv-webgpu-prepare \
+// RUN: mlir-opt %s -test-vulkan-runner-pipeline=spirv-webgpu-prepare \
// RUN: | mlir-vulkan-runner - \
// RUN: --shared-libs=%vulkan-runtime-wrappers,%mlir_runner_utils \
// RUN: --entry-point-result=void | FileCheck %s
diff --git a/mlir/test/mlir-vulkan-runner/umul_extended.mlir b/mlir/test/mlir-vulkan-runner/umul_extended.mlir
index 78300d2fd81dd5..803b8c3d336d33 100644
--- a/mlir/test/mlir-vulkan-runner/umul_extended.mlir
+++ b/mlir/test/mlir-vulkan-runner/umul_extended.mlir
@@ -6,7 +6,7 @@
// RUN: --shared-libs=%vulkan-runtime-wrappers,%mlir_runner_utils \
// RUN: --entry-point-result=void | FileCheck %s
-// RUN: mlir-opt %s -test-vulkan-runner-pipeline -spirv-webgpu-prepare \
+// RUN: mlir-opt %s -test-vulkan-runner-pipeline=spirv-webgpu-prepare \
// RUN: | mlir-vulkan-runner - \
// RUN: --shared-libs=%vulkan-runtime-wrappers,%mlir_runner_utils \
// RUN: --entry-point-result=void | FileCheck %s
|
andfau-amd
force-pushed
the
73457-runner-migration-vulkan-pipeline-with-serialization
branch
2 times, most recently
from
d7256b9 to
f64d825
Compare
kuhar
reviewed
Jan 2, 2025
mlir/lib/Conversion/GPUToVulkan/ConvertGPULaunchFuncToVulkanLaunchFunc.cpp
Outdated
Show resolved
Hide resolved
mlir/lib/Conversion/GPUToVulkan/ConvertGPULaunchFuncToVulkanLaunchFunc.cpp
Outdated
Show resolved
Hide resolved
mlir/test/Conversion/ConvertToSPIRV/convert-gpu-modules-nested.mlir
Outdated
Show resolved
Hide resolved
andfau-amd
force-pushed
the
73457-runner-migration-vulkan-pipeline-with-serialization
branch
from
f64d825 to
10eed7b
Compare
kuhar
reviewed
Jan 2, 2025
mlir/lib/Conversion/GPUToVulkan/ConvertGPULaunchFuncToVulkanLaunchFunc.cpp
Outdated
Show resolved
Hide resolved
This commit is a further incremental step toward moving the whole mlir-vulkan-runner MLIR pass pipeline into mlir-opt (see llvm#73457). The previous step was b225b3adf7b78387c9fcb97a3ff0e0a1e26eafe2, which moved all device passes prior to SPIR-V serialization into a new mlir-opt test pass, `-test-vulkan-runner-pipeline`. This commit changes how SPIR-V serialization is accomplished for Vulkan runner tests. Until now, this was done by the Vulkan-specific ConvertGpuLaunchFuncToVulkanLaunchFunc pass. With this commit, this responsibility is removed from that pass, and is instead done with the existing generic GpuModuleToBinaryPass. In addition, the SPIR-V serialization step is no longer done inside mlir-vulkan-runner, but rather inside mlir-opt (in the `-test-vulkan-runner-pipeline` pass). Both of these changes represent a greater alignment between mlir-vulkan-runner and the other GPU integration tests. Notably, the IR shapes produced by the mlir-opt pipelines for the Vulkan and SYCL runners are now much more similar, with both using a gpu.binary op for the serialized SPIR-V kernel. In order to enable this, this commit includes these supporting changes: - ConvertToSPIRVPass is enhanced to support producing the IR shape where a spirv.module is nested inside a gpu.module, since this is what GpuModuleToBinaryPass expects. - ConvertGPULaunchFuncToVulkanLaunchFunc is changed to remove its SPIR-V serialization functionality, and instead now extracts the SPIR-V from a gpu.binary operation (as produced by ConvertToSPIRVPass). - `-test-vulkan-runner-pipeline` now attaches SPIR-V target information required by GpuModuleToBinaryPass. - The WebGPU pass option, which had been removed from mlir-vulkan-runner in the previous commit in this series, is restored as an option to `-test-vulkan-runner-pipeline` instead, so that the WebGPU pass continues being inserted into the pipeline just before SPIR-V serialization.
andfau-amd
force-pushed
the
73457-runner-migration-vulkan-pipeline-with-serialization
branch
from
10eed7b to
6764075
Compare
Hardcode84
approved these changes
Jan 9, 2025
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 is a further incremental step toward moving the whole mlir-vulkan-runner MLIR pass pipeline into mlir-opt (see #73457). The previous step was b225b3adf7b78387c9fcb97a3ff0e0a1e26eafe2, which moved all device passes prior to SPIR-V serialization into a new mlir-opt test pass,
-test-vulkan-runner-pipeline.This commit changes how SPIR-V serialization is accomplished for Vulkan runner tests. Until now, this was done by the Vulkan-specific ConvertGpuLaunchFuncToVulkanLaunchFunc pass. With this commit, this responsibility is removed from that pass, and is instead done with the existing generic GpuModuleToBinaryPass. In addition, the SPIR-V serialization step is no longer done inside mlir-vulkan-runner, but rather inside mlir-opt (in the
-test-vulkan-runner-pipelinepass). Both of these changes represent a greater alignment between mlir-vulkan-runner and the other GPU integration tests. Notably, the IR shapes produced by the mlir-opt pipelines for the Vulkan and SYCL runners are now much more similar, with both using a gpu.binary op for the serialized SPIR-V kernel.In order to enable this, this commit includes these supporting changes:
-test-vulkan-runner-pipelinenow attaches SPIR-V target information required by GpuModuleToBinaryPass.-test-vulkan-runner-pipelineinstead, so that the WebGPU pass continues being inserted into the pipeline just before SPIR-V serialization.