[Mlir-commits] [mlir] [mlir] Remove mlir-vulkan-runner and GPUToVulkan conversion passes (PR #123750)
Andrea Faulds
llvmlistbot at llvm.org
Tue Jan 21 05:55:20 PST 2025
https://github.com/andfau-amd created https://github.com/llvm/llvm-project/pull/123750
This follows up on 733be4ed7dcf976719f424c0cb81b77a14f91f5a, which made mlir-vulkan-runner and its associated passes redundant, and completes the main goal of #73457. The mlir-vulkan-runner tests become part of the integration test suite, and the Vulkan runner runtime components become part of ExecutionEngine, just as was done when removing other target-specific runners.
>From 9dc6e2e04b055a212b2e045bcef45e17d944401f Mon Sep 17 00:00:00 2001
From: Andrea Faulds <andrea.faulds at amd.com>
Date: Tue, 21 Jan 2025 14:48:52 +0100
Subject: [PATCH] [mlir] Remove mlir-vulkan-runner and GPUToVulkan conversion
passes
This follows up on 733be4ed7dcf976719f424c0cb81b77a14f91f5a, which made
mlir-vulkan-runner and its associated passes redundant, and completes
the main goal of #73457. The mlir-vulkan-runner tests become part of the
integration test suite, and the Vulkan runner runtime components become
part of ExecutionEngine, just as was done when removing other
target-specific runners.
---
.../GPUToVulkan/ConvertGPUToVulkanPass.h | 36 --
mlir/include/mlir/Conversion/Passes.h | 1 -
mlir/include/mlir/Conversion/Passes.td | 25 -
mlir/lib/Conversion/CMakeLists.txt | 1 -
.../lib/Conversion/GPUToVulkan/CMakeLists.txt | 19 -
...ConvertGPULaunchFuncToVulkanLaunchFunc.cpp | 219 ---------
.../ConvertLaunchFuncToVulkanCalls.cpp | 448 ------------------
mlir/lib/ExecutionEngine/CMakeLists.txt | 45 ++
.../ExecutionEngine}/VulkanRuntime.cpp | 0
.../ExecutionEngine}/VulkanRuntime.h | 0
.../VulkanRuntimeWrappers.cpp} | 77 +--
mlir/test/CMakeLists.txt | 2 +-
.../Conversion/GPUToVulkan/invoke-vulkan.mlir | 62 ---
.../lower-gpu-launch-vulkan-launch.mlir | 35 --
.../GPU/Vulkan}/addf.mlir | 4 +-
.../GPU/Vulkan}/addf_if.mlir | 4 +-
.../GPU/Vulkan}/addi.mlir | 4 +-
.../GPU/Vulkan}/addi8.mlir | 4 +-
.../GPU/Vulkan}/addui_extended.mlir | 8 +-
.../GPU/Vulkan}/lit.local.cfg | 0
.../GPU/Vulkan}/mulf.mlir | 4 +-
.../GPU/Vulkan}/smul_extended.mlir | 8 +-
.../GPU/Vulkan}/subf.mlir | 4 +-
.../GPU/Vulkan}/time.mlir | 4 +-
.../GPU/Vulkan}/umul_extended.mlir | 8 +-
.../GPU/Vulkan}/vector-deinterleave.mlir | 4 +-
.../GPU/Vulkan}/vector-interleave.mlir | 4 +-
.../GPU/Vulkan}/vector-shuffle.mlir | 4 +-
.../lib/Pass/TestVulkanRunnerPipeline.cpp | 29 +-
mlir/test/lit.cfg.py | 2 +-
mlir/tools/CMakeLists.txt | 1 -
mlir/tools/mlir-vulkan-runner/CMakeLists.txt | 99 ----
.../mlir-vulkan-runner/mlir-vulkan-runner.cpp | 88 ----
33 files changed, 95 insertions(+), 1158 deletions(-)
delete mode 100644 mlir/include/mlir/Conversion/GPUToVulkan/ConvertGPUToVulkanPass.h
delete mode 100644 mlir/lib/Conversion/GPUToVulkan/CMakeLists.txt
delete mode 100644 mlir/lib/Conversion/GPUToVulkan/ConvertGPULaunchFuncToVulkanLaunchFunc.cpp
delete mode 100644 mlir/lib/Conversion/GPUToVulkan/ConvertLaunchFuncToVulkanCalls.cpp
rename mlir/{tools/mlir-vulkan-runner => lib/ExecutionEngine}/VulkanRuntime.cpp (100%)
rename mlir/{tools/mlir-vulkan-runner => lib/ExecutionEngine}/VulkanRuntime.h (100%)
rename mlir/{tools/mlir-vulkan-runner/vulkan-runtime-wrappers.cpp => lib/ExecutionEngine/VulkanRuntimeWrappers.cpp} (75%)
delete mode 100644 mlir/test/Conversion/GPUToVulkan/invoke-vulkan.mlir
delete mode 100644 mlir/test/Conversion/GPUToVulkan/lower-gpu-launch-vulkan-launch.mlir
rename mlir/test/{mlir-vulkan-runner => Integration/GPU/Vulkan}/addf.mlir (91%)
rename mlir/test/{mlir-vulkan-runner => Integration/GPU/Vulkan}/addf_if.mlir (91%)
rename mlir/test/{mlir-vulkan-runner => Integration/GPU/Vulkan}/addi.mlir (91%)
rename mlir/test/{mlir-vulkan-runner => Integration/GPU/Vulkan}/addi8.mlir (92%)
rename mlir/test/{mlir-vulkan-runner => Integration/GPU/Vulkan}/addui_extended.mlir (90%)
rename mlir/test/{mlir-vulkan-runner => Integration/GPU/Vulkan}/lit.local.cfg (100%)
rename mlir/test/{mlir-vulkan-runner => Integration/GPU/Vulkan}/mulf.mlir (91%)
rename mlir/test/{mlir-vulkan-runner => Integration/GPU/Vulkan}/smul_extended.mlir (90%)
rename mlir/test/{mlir-vulkan-runner => Integration/GPU/Vulkan}/subf.mlir (92%)
rename mlir/test/{mlir-vulkan-runner => Integration/GPU/Vulkan}/time.mlir (92%)
rename mlir/test/{mlir-vulkan-runner => Integration/GPU/Vulkan}/umul_extended.mlir (90%)
rename mlir/test/{mlir-vulkan-runner => Integration/GPU/Vulkan}/vector-deinterleave.mlir (96%)
rename mlir/test/{mlir-vulkan-runner => Integration/GPU/Vulkan}/vector-interleave.mlir (96%)
rename mlir/test/{mlir-vulkan-runner => Integration/GPU/Vulkan}/vector-shuffle.mlir (96%)
delete mode 100644 mlir/tools/mlir-vulkan-runner/CMakeLists.txt
delete mode 100644 mlir/tools/mlir-vulkan-runner/mlir-vulkan-runner.cpp
diff --git a/mlir/include/mlir/Conversion/GPUToVulkan/ConvertGPUToVulkanPass.h b/mlir/include/mlir/Conversion/GPUToVulkan/ConvertGPUToVulkanPass.h
deleted file mode 100644
index f69720328f2a42..00000000000000
--- a/mlir/include/mlir/Conversion/GPUToVulkan/ConvertGPUToVulkanPass.h
+++ /dev/null
@@ -1,36 +0,0 @@
-//===- ConvertGPUToVulkanPass.h - GPU to Vulkan conversion pass -*- C++ -*-===//
-//
-// 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
-//
-//===----------------------------------------------------------------------===//
-//
-// The file declares a pass to convert GPU dialect ops to to Vulkan runtime
-// calls.
-//
-//===----------------------------------------------------------------------===//
-
-#ifndef MLIR_CONVERSION_GPUTOVULKAN_CONVERTGPUTOVULKANPASS_H
-#define MLIR_CONVERSION_GPUTOVULKAN_CONVERTGPUTOVULKANPASS_H
-
-#include "mlir/Support/LLVM.h"
-
-#include <memory>
-
-namespace mlir {
-
-class ModuleOp;
-template <typename T>
-class OperationPass;
-class Pass;
-
-#define GEN_PASS_DECL_CONVERTVULKANLAUNCHFUNCTOVULKANCALLSPASS
-#define GEN_PASS_DECL_CONVERTGPULAUNCHFUNCTOVULKANLAUNCHFUNC
-#include "mlir/Conversion/Passes.h.inc"
-
-std::unique_ptr<OperationPass<mlir::ModuleOp>>
-createConvertGpuLaunchFuncToVulkanLaunchFuncPass();
-
-} // namespace mlir
-#endif // MLIR_CONVERSION_GPUTOVULKAN_CONVERTGPUTOVULKANPASS_H
diff --git a/mlir/include/mlir/Conversion/Passes.h b/mlir/include/mlir/Conversion/Passes.h
index 6a564e9bfc5a9b..e9761c20642c0f 100644
--- a/mlir/include/mlir/Conversion/Passes.h
+++ b/mlir/include/mlir/Conversion/Passes.h
@@ -39,7 +39,6 @@
#include "mlir/Conversion/GPUToNVVM/GPUToNVVMPass.h"
#include "mlir/Conversion/GPUToROCDL/GPUToROCDLPass.h"
#include "mlir/Conversion/GPUToSPIRV/GPUToSPIRVPass.h"
-#include "mlir/Conversion/GPUToVulkan/ConvertGPUToVulkanPass.h"
#include "mlir/Conversion/IndexToLLVM/IndexToLLVM.h"
#include "mlir/Conversion/IndexToSPIRV/IndexToSPIRV.h"
#include "mlir/Conversion/LinalgToStandard/LinalgToStandard.h"
diff --git a/mlir/include/mlir/Conversion/Passes.td b/mlir/include/mlir/Conversion/Passes.td
index ac417fff27eb80..f9fa2d4595f999 100644
--- a/mlir/include/mlir/Conversion/Passes.td
+++ b/mlir/include/mlir/Conversion/Passes.td
@@ -654,31 +654,6 @@ def ConvertGPUToSPIRV : Pass<"convert-gpu-to-spirv", "ModuleOp"> {
];
}
-//===----------------------------------------------------------------------===//
-// GPUToVulkan
-//===----------------------------------------------------------------------===//
-
-def ConvertGpuLaunchFuncToVulkanLaunchFunc
- : Pass<"convert-gpu-launch-to-vulkan-launch", "ModuleOp"> {
- let summary = "Convert gpu.launch_func to vulkanLaunch external call";
- let description = [{
- This pass is only intended for the mlir-vulkan-runner.
- }];
- let constructor = "mlir::createConvertGpuLaunchFuncToVulkanLaunchFuncPass()";
- let dependentDialects = ["spirv::SPIRVDialect"];
-}
-
-def ConvertVulkanLaunchFuncToVulkanCallsPass
- : Pass<"launch-func-to-vulkan", "ModuleOp"> {
- let summary = "Convert vulkanLaunch external call to Vulkan runtime external "
- "calls";
- let description = [{
- This pass is only intended for the mlir-vulkan-runner.
- }];
-
- let dependentDialects = ["LLVM::LLVMDialect"];
-}
-
//===----------------------------------------------------------------------===//
// ConvertIndexToLLVMPass
//===----------------------------------------------------------------------===//
diff --git a/mlir/lib/Conversion/CMakeLists.txt b/mlir/lib/Conversion/CMakeLists.txt
index 791e94e491587c..a570978f03757c 100644
--- a/mlir/lib/Conversion/CMakeLists.txt
+++ b/mlir/lib/Conversion/CMakeLists.txt
@@ -28,7 +28,6 @@ add_subdirectory(GPUToLLVMSPV)
add_subdirectory(GPUToNVVM)
add_subdirectory(GPUToROCDL)
add_subdirectory(GPUToSPIRV)
-add_subdirectory(GPUToVulkan)
add_subdirectory(IndexToLLVM)
add_subdirectory(IndexToSPIRV)
add_subdirectory(LinalgToStandard)
diff --git a/mlir/lib/Conversion/GPUToVulkan/CMakeLists.txt b/mlir/lib/Conversion/GPUToVulkan/CMakeLists.txt
deleted file mode 100644
index faeb32f2bc8cd3..00000000000000
--- a/mlir/lib/Conversion/GPUToVulkan/CMakeLists.txt
+++ /dev/null
@@ -1,19 +0,0 @@
-add_mlir_conversion_library(MLIRGPUToVulkanTransforms
- ConvertLaunchFuncToVulkanCalls.cpp
- ConvertGPULaunchFuncToVulkanLaunchFunc.cpp
-
- DEPENDS
- MLIRConversionPassIncGen
-
- LINK_LIBS PUBLIC
- MLIRFuncDialect
- MLIRGPUDialect
- MLIRIR
- MLIRLLVMDialect
- MLIRPass
- MLIRSPIRVDialect
- MLIRSPIRVSerialization
- MLIRSupport
- MLIRTransforms
- MLIRTranslateLib
- )
diff --git a/mlir/lib/Conversion/GPUToVulkan/ConvertGPULaunchFuncToVulkanLaunchFunc.cpp b/mlir/lib/Conversion/GPUToVulkan/ConvertGPULaunchFuncToVulkanLaunchFunc.cpp
deleted file mode 100644
index 8488fac69e8e31..00000000000000
--- a/mlir/lib/Conversion/GPUToVulkan/ConvertGPULaunchFuncToVulkanLaunchFunc.cpp
+++ /dev/null
@@ -1,219 +0,0 @@
-//===- ConvertGPULaunchFuncToVulkanLaunchFunc.cpp - MLIR conversion pass --===//
-//
-// 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 to convert gpu launch function into a vulkan
-// 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.
-//
-//===----------------------------------------------------------------------===//
-
-#include "mlir/Conversion/GPUToVulkan/ConvertGPUToVulkanPass.h"
-
-#include "mlir/Dialect/Func/IR/FuncOps.h"
-#include "mlir/Dialect/GPU/IR/GPUDialect.h"
-#include "mlir/Dialect/SPIRV/IR/SPIRVDialect.h"
-#include "mlir/Dialect/SPIRV/IR/SPIRVOps.h"
-#include "mlir/IR/Attributes.h"
-#include "mlir/IR/Builders.h"
-#include "mlir/IR/BuiltinOps.h"
-#include "mlir/IR/BuiltinTypes.h"
-#include "mlir/Pass/Pass.h"
-#include "mlir/Target/SPIRV/Serialization.h"
-
-namespace mlir {
-#define GEN_PASS_DEF_CONVERTGPULAUNCHFUNCTOVULKANLAUNCHFUNC
-#include "mlir/Conversion/Passes.h.inc"
-} // namespace mlir
-
-using namespace mlir;
-
-static constexpr const char *kSPIRVBlobAttrName = "spirv_blob";
-static constexpr const char *kSPIRVEntryPointAttrName = "spirv_entry_point";
-static constexpr const char *kSPIRVElementTypesAttrName = "spirv_element_types";
-static constexpr const char *kVulkanLaunch = "vulkanLaunch";
-
-namespace {
-
-/// 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> {
-public:
- void runOnOperation() override;
-
-private:
- /// Extracts a SPIR-V binary shader from the given `module`, if any.
- /// Note that this also removes the binary from the IR.
- FailureOr<StringAttr> getBinaryShader(ModuleOp module);
-
- /// Converts the given `launchOp` to vulkan launch call.
- void convertGpuLaunchFunc(gpu::LaunchFuncOp launchOp);
-
- /// Checks where the given type is supported by Vulkan runtime.
- bool isSupportedType(Type type) {
- if (auto memRefType = dyn_cast_or_null<MemRefType>(type)) {
- auto elementType = memRefType.getElementType();
- return memRefType.hasRank() &&
- (memRefType.getRank() >= 1 && memRefType.getRank() <= 3) &&
- (elementType.isIntOrFloat());
- }
- return false;
- }
-
- /// Declares the vulkan launch function. Returns an error if the any type of
- /// operand is unsupported by Vulkan runtime.
- LogicalResult declareVulkanLaunchFunc(Location loc,
- gpu::LaunchFuncOp launchOp);
-
-private:
- /// The number of vulkan launch configuration operands, placed at the leading
- /// positions of the operand list.
- static constexpr unsigned kVulkanLaunchNumConfigOperands = 3;
-};
-
-} // namespace
-
-void ConvertGpuLaunchFuncToVulkanLaunchFunc::runOnOperation() {
- bool done = false;
- getOperation().walk([this, &done](gpu::LaunchFuncOp op) {
- if (done) {
- op.emitError("should only contain one 'gpu::LaunchFuncOp' op");
- return signalPassFailure();
- }
- done = true;
- convertGpuLaunchFunc(op);
- });
-
- // Erase `gpu::GPUModuleOp` and `spirv::Module` operations.
- for (auto gpuModule :
- llvm::make_early_inc_range(getOperation().getOps<gpu::GPUModuleOp>()))
- gpuModule.erase();
-
- for (auto spirvModule :
- llvm::make_early_inc_range(getOperation().getOps<spirv::ModuleOp>()))
- spirvModule.erase();
-}
-
-LogicalResult ConvertGpuLaunchFuncToVulkanLaunchFunc::declareVulkanLaunchFunc(
- Location loc, gpu::LaunchFuncOp launchOp) {
- auto builder = OpBuilder::atBlockEnd(getOperation().getBody());
-
- // Workgroup size is written into the kernel. So to properly modelling
- // vulkan launch, we have to skip local workgroup size configuration here.
- SmallVector<Type, 8> gpuLaunchTypes(launchOp.getOperandTypes());
- // The first kVulkanLaunchNumConfigOperands of the gpu.launch_func op are the
- // same as the config operands for the vulkan launch call op.
- SmallVector<Type, 8> vulkanLaunchTypes(gpuLaunchTypes.begin(),
- gpuLaunchTypes.begin() +
- kVulkanLaunchNumConfigOperands);
- vulkanLaunchTypes.append(gpuLaunchTypes.begin() +
- gpu::LaunchOp::kNumConfigOperands,
- gpuLaunchTypes.end());
-
- // Check that all operands have supported types except those for the
- // launch configuration.
- for (auto type :
- llvm::drop_begin(vulkanLaunchTypes, kVulkanLaunchNumConfigOperands)) {
- if (!isSupportedType(type))
- return launchOp.emitError() << type << " is unsupported to run on Vulkan";
- }
-
- // Declare vulkan launch function.
- auto funcType = builder.getFunctionType(vulkanLaunchTypes, {});
- builder.create<func::FuncOp>(loc, kVulkanLaunch, funcType).setPrivate();
-
- return success();
-}
-
-FailureOr<StringAttr>
-ConvertGpuLaunchFuncToVulkanLaunchFunc::getBinaryShader(ModuleOp module) {
- bool done = false;
- StringAttr binaryAttr;
- gpu::BinaryOp binaryToErase;
- for (auto gpuBinary : module.getOps<gpu::BinaryOp>()) {
- if (done)
- return gpuBinary.emitError("should only contain one 'gpu.binary' op");
- done = true;
-
- 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");
-
- binaryAttr = object.getObject();
- binaryToErase = gpuBinary;
- }
- if (!done)
- return module.emitError("should contain a 'gpu.binary' op");
-
- // Remove the binary to avoid confusing later conversion passes.
- binaryToErase.erase();
- return binaryAttr;
-}
-
-void ConvertGpuLaunchFuncToVulkanLaunchFunc::convertGpuLaunchFunc(
- gpu::LaunchFuncOp launchOp) {
- ModuleOp module = getOperation();
- OpBuilder builder(launchOp);
- Location loc = launchOp.getLoc();
-
- FailureOr<StringAttr> binaryAttr = getBinaryShader(module);
- // Extract SPIR-V from `gpu.binary` op.
- if (failed(binaryAttr))
- return signalPassFailure();
-
- // Declare vulkan launch function.
- if (failed(declareVulkanLaunchFunc(loc, launchOp)))
- return signalPassFailure();
-
- SmallVector<Value, 8> gpuLaunchOperands(launchOp.getOperands());
- SmallVector<Value, 8> vulkanLaunchOperands(
- gpuLaunchOperands.begin(),
- gpuLaunchOperands.begin() + kVulkanLaunchNumConfigOperands);
- vulkanLaunchOperands.append(gpuLaunchOperands.begin() +
- gpu::LaunchOp::kNumConfigOperands,
- gpuLaunchOperands.end());
-
- // Create vulkan launch call op.
- auto vulkanLaunchCallOp = builder.create<func::CallOp>(
- loc, TypeRange{}, SymbolRefAttr::get(builder.getContext(), kVulkanLaunch),
- vulkanLaunchOperands);
-
- // Set SPIR-V binary shader data as an attribute.
- vulkanLaunchCallOp->setAttr(kSPIRVBlobAttrName, *binaryAttr);
-
- // Set entry point name as an attribute.
- vulkanLaunchCallOp->setAttr(kSPIRVEntryPointAttrName,
- launchOp.getKernelName());
-
- // Add MemRef element types before they're lost when lowering to LLVM.
- SmallVector<Type> elementTypes;
- for (Type type : llvm::drop_begin(launchOp.getOperandTypes(),
- gpu::LaunchOp::kNumConfigOperands)) {
- // The below cast always succeeds as it has already been verified in
- // 'declareVulkanLaunchFunc' that these are MemRefs with compatible element
- // types.
- elementTypes.push_back(cast<MemRefType>(type).getElementType());
- }
- vulkanLaunchCallOp->setAttr(kSPIRVElementTypesAttrName,
- builder.getTypeArrayAttr(elementTypes));
-
- launchOp.erase();
-}
-
-std::unique_ptr<mlir::OperationPass<mlir::ModuleOp>>
-mlir::createConvertGpuLaunchFuncToVulkanLaunchFuncPass() {
- return std::make_unique<ConvertGpuLaunchFuncToVulkanLaunchFunc>();
-}
diff --git a/mlir/lib/Conversion/GPUToVulkan/ConvertLaunchFuncToVulkanCalls.cpp b/mlir/lib/Conversion/GPUToVulkan/ConvertLaunchFuncToVulkanCalls.cpp
deleted file mode 100644
index 938db549630680..00000000000000
--- a/mlir/lib/Conversion/GPUToVulkan/ConvertLaunchFuncToVulkanCalls.cpp
+++ /dev/null
@@ -1,448 +0,0 @@
-//===- ConvertLaunchFuncToVulkanCalls.cpp - MLIR Vulkan conversion passes -===//
-//
-// 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 to convert vulkan launch call into a sequence of
-// Vulkan runtime calls. The Vulkan runtime API surface is huge so currently we
-// don't expose separate external functions in IR for each of them, instead we
-// expose a few external functions to wrapper libraries which manages Vulkan
-// runtime.
-//
-//===----------------------------------------------------------------------===//
-
-#include "mlir/Conversion/GPUToVulkan/ConvertGPUToVulkanPass.h"
-
-#include "mlir/Dialect/LLVMIR/LLVMDialect.h"
-#include "mlir/IR/Attributes.h"
-#include "mlir/IR/Builders.h"
-#include "mlir/IR/BuiltinOps.h"
-#include "mlir/Pass/Pass.h"
-#include "llvm/ADT/SmallString.h"
-#include "llvm/Support/FormatVariadic.h"
-
-namespace mlir {
-#define GEN_PASS_DEF_CONVERTVULKANLAUNCHFUNCTOVULKANCALLSPASS
-#include "mlir/Conversion/Passes.h.inc"
-} // namespace mlir
-
-using namespace mlir;
-
-static constexpr const char *kCInterfaceVulkanLaunch =
- "_mlir_ciface_vulkanLaunch";
-static constexpr const char *kDeinitVulkan = "deinitVulkan";
-static constexpr const char *kRunOnVulkan = "runOnVulkan";
-static constexpr const char *kInitVulkan = "initVulkan";
-static constexpr const char *kSetBinaryShader = "setBinaryShader";
-static constexpr const char *kSetEntryPoint = "setEntryPoint";
-static constexpr const char *kSetNumWorkGroups = "setNumWorkGroups";
-static constexpr const char *kSPIRVBinary = "SPIRV_BIN";
-static constexpr const char *kSPIRVBlobAttrName = "spirv_blob";
-static constexpr const char *kSPIRVEntryPointAttrName = "spirv_entry_point";
-static constexpr const char *kSPIRVElementTypesAttrName = "spirv_element_types";
-static constexpr const char *kVulkanLaunch = "vulkanLaunch";
-
-namespace {
-
-/// A pass to convert vulkan launch call op into a sequence of Vulkan
-/// runtime calls in the following order:
-///
-/// * initVulkan -- initializes vulkan runtime
-/// * bindMemRef -- binds memref
-/// * setBinaryShader -- sets the binary shader data
-/// * setEntryPoint -- sets the entry point name
-/// * setNumWorkGroups -- sets the number of a local workgroups
-/// * runOnVulkan -- runs vulkan runtime
-/// * deinitVulkan -- deinitializes vulkan runtime
-///
-class VulkanLaunchFuncToVulkanCallsPass
- : public impl::ConvertVulkanLaunchFuncToVulkanCallsPassBase<
- VulkanLaunchFuncToVulkanCallsPass> {
-private:
- void initializeCachedTypes() {
- llvmFloatType = Float32Type::get(&getContext());
- llvmVoidType = LLVM::LLVMVoidType::get(&getContext());
- llvmPointerType = LLVM::LLVMPointerType::get(&getContext());
- llvmInt32Type = IntegerType::get(&getContext(), 32);
- llvmInt64Type = IntegerType::get(&getContext(), 64);
- }
-
- Type getMemRefType(uint32_t rank, Type elemenType) {
- // According to the MLIR doc memref argument is converted into a
- // pointer-to-struct argument of type:
- // template <typename Elem, size_t Rank>
- // struct {
- // Elem *allocated;
- // Elem *aligned;
- // int64_t offset;
- // int64_t sizes[Rank]; // omitted when rank == 0
- // int64_t strides[Rank]; // omitted when rank == 0
- // };
- auto llvmArrayRankElementSizeType =
- LLVM::LLVMArrayType::get(getInt64Type(), rank);
-
- // Create a type
- // `!llvm<"{ `element-type`*, `element-type`*, i64,
- // [`rank` x i64], [`rank` x i64]}">`.
- return LLVM::LLVMStructType::getLiteral(
- &getContext(),
- {llvmPointerType, llvmPointerType, getInt64Type(),
- llvmArrayRankElementSizeType, llvmArrayRankElementSizeType});
- }
-
- Type getVoidType() { return llvmVoidType; }
- Type getPointerType() { return llvmPointerType; }
- Type getInt32Type() { return llvmInt32Type; }
- Type getInt64Type() { return llvmInt64Type; }
-
- /// Creates an LLVM global for the given `name`.
- Value createEntryPointNameConstant(StringRef name, Location loc,
- OpBuilder &builder);
-
- /// Declares all needed runtime functions.
- void declareVulkanFunctions(Location loc);
-
- /// Checks whether the given LLVM::CallOp is a vulkan launch call op.
- bool isVulkanLaunchCallOp(LLVM::CallOp callOp) {
- return (callOp.getCallee() && *callOp.getCallee() == kVulkanLaunch &&
- callOp.getNumOperands() >= kVulkanLaunchNumConfigOperands);
- }
-
- /// Checks whether the given LLVM::CallOp is a "ci_face" vulkan launch call
- /// op.
- bool isCInterfaceVulkanLaunchCallOp(LLVM::CallOp callOp) {
- return (callOp.getCallee() &&
- *callOp.getCallee() == kCInterfaceVulkanLaunch &&
- callOp.getNumOperands() >= kVulkanLaunchNumConfigOperands);
- }
-
- /// Translates the given `vulkanLaunchCallOp` to the sequence of Vulkan
- /// runtime calls.
- void translateVulkanLaunchCall(LLVM::CallOp vulkanLaunchCallOp);
-
- /// Creates call to `bindMemRef` for each memref operand.
- void createBindMemRefCalls(LLVM::CallOp vulkanLaunchCallOp,
- Value vulkanRuntime);
-
- /// Collects SPIRV attributes from the given `vulkanLaunchCallOp`.
- void collectSPIRVAttributes(LLVM::CallOp vulkanLaunchCallOp);
-
- /// Deduces a rank from the given 'launchCallArg`.
- LogicalResult deduceMemRefRank(Value launchCallArg, uint32_t &rank);
-
- /// Returns a string representation from the given `type`.
- StringRef stringifyType(Type type) {
- if (isa<Float32Type>(type))
- return "Float";
- if (isa<Float16Type>(type))
- return "Half";
- if (auto intType = dyn_cast<IntegerType>(type)) {
- if (intType.getWidth() == 32)
- return "Int32";
- if (intType.getWidth() == 16)
- return "Int16";
- if (intType.getWidth() == 8)
- return "Int8";
- }
-
- llvm_unreachable("unsupported type");
- }
-
-public:
- using Base::Base;
-
- void runOnOperation() override;
-
-private:
- Type llvmFloatType;
- Type llvmVoidType;
- Type llvmPointerType;
- Type llvmInt32Type;
- Type llvmInt64Type;
-
- struct SPIRVAttributes {
- StringAttr blob;
- StringAttr entryPoint;
- SmallVector<Type> elementTypes;
- };
-
- // TODO: Use an associative array to support multiple vulkan launch calls.
- SPIRVAttributes spirvAttributes;
- /// The number of vulkan launch configuration operands, placed at the leading
- /// positions of the operand list.
- static constexpr unsigned kVulkanLaunchNumConfigOperands = 3;
-};
-
-} // namespace
-
-void VulkanLaunchFuncToVulkanCallsPass::runOnOperation() {
- initializeCachedTypes();
-
- // Collect SPIR-V attributes such as `spirv_blob` and
- // `spirv_entry_point_name`.
- getOperation().walk([this](LLVM::CallOp op) {
- if (isVulkanLaunchCallOp(op))
- collectSPIRVAttributes(op);
- });
-
- // Convert vulkan launch call op into a sequence of Vulkan runtime calls.
- getOperation().walk([this](LLVM::CallOp op) {
- if (isCInterfaceVulkanLaunchCallOp(op))
- translateVulkanLaunchCall(op);
- });
-}
-
-void VulkanLaunchFuncToVulkanCallsPass::collectSPIRVAttributes(
- LLVM::CallOp vulkanLaunchCallOp) {
- // Check that `kSPIRVBinary` and `kSPIRVEntryPoint` are present in attributes
- // for the given vulkan launch call.
- auto spirvBlobAttr =
- vulkanLaunchCallOp->getAttrOfType<StringAttr>(kSPIRVBlobAttrName);
- if (!spirvBlobAttr) {
- vulkanLaunchCallOp.emitError()
- << "missing " << kSPIRVBlobAttrName << " attribute";
- return signalPassFailure();
- }
-
- auto spirvEntryPointNameAttr =
- vulkanLaunchCallOp->getAttrOfType<StringAttr>(kSPIRVEntryPointAttrName);
- if (!spirvEntryPointNameAttr) {
- vulkanLaunchCallOp.emitError()
- << "missing " << kSPIRVEntryPointAttrName << " attribute";
- return signalPassFailure();
- }
-
- auto spirvElementTypesAttr =
- vulkanLaunchCallOp->getAttrOfType<ArrayAttr>(kSPIRVElementTypesAttrName);
- if (!spirvElementTypesAttr) {
- vulkanLaunchCallOp.emitError()
- << "missing " << kSPIRVElementTypesAttrName << " attribute";
- return signalPassFailure();
- }
- if (llvm::any_of(spirvElementTypesAttr,
- [](Attribute attr) { return !isa<TypeAttr>(attr); })) {
- vulkanLaunchCallOp.emitError()
- << "expected " << spirvElementTypesAttr << " to be an array of types";
- return signalPassFailure();
- }
-
- spirvAttributes.blob = spirvBlobAttr;
- spirvAttributes.entryPoint = spirvEntryPointNameAttr;
- spirvAttributes.elementTypes =
- llvm::to_vector(spirvElementTypesAttr.getAsValueRange<mlir::TypeAttr>());
-}
-
-void VulkanLaunchFuncToVulkanCallsPass::createBindMemRefCalls(
- LLVM::CallOp cInterfaceVulkanLaunchCallOp, Value vulkanRuntime) {
- if (cInterfaceVulkanLaunchCallOp.getNumOperands() ==
- kVulkanLaunchNumConfigOperands)
- return;
- OpBuilder builder(cInterfaceVulkanLaunchCallOp);
- Location loc = cInterfaceVulkanLaunchCallOp.getLoc();
-
- // Create LLVM constant for the descriptor set index.
- // Bind all memrefs to the `0` descriptor set, the same way as `GPUToSPIRV`
- // pass does.
- Value descriptorSet =
- builder.create<LLVM::ConstantOp>(loc, getInt32Type(), 0);
-
- for (auto [index, ptrToMemRefDescriptor] :
- llvm::enumerate(cInterfaceVulkanLaunchCallOp.getOperands().drop_front(
- kVulkanLaunchNumConfigOperands))) {
- // Create LLVM constant for the descriptor binding index.
- Value descriptorBinding =
- builder.create<LLVM::ConstantOp>(loc, getInt32Type(), index);
-
- if (index >= spirvAttributes.elementTypes.size()) {
- cInterfaceVulkanLaunchCallOp.emitError()
- << kSPIRVElementTypesAttrName << " missing element type for "
- << ptrToMemRefDescriptor;
- return signalPassFailure();
- }
-
- uint32_t rank = 0;
- Type type = spirvAttributes.elementTypes[index];
- if (failed(deduceMemRefRank(ptrToMemRefDescriptor, rank))) {
- cInterfaceVulkanLaunchCallOp.emitError()
- << "invalid memref descriptor " << ptrToMemRefDescriptor.getType();
- return signalPassFailure();
- }
-
- auto symbolName =
- llvm::formatv("bindMemRef{0}D{1}", rank, stringifyType(type)).str();
- // Create call to `bindMemRef`.
- builder.create<LLVM::CallOp>(
- loc, TypeRange(), StringRef(symbolName.data(), symbolName.size()),
- ValueRange{vulkanRuntime, descriptorSet, descriptorBinding,
- ptrToMemRefDescriptor});
- }
-}
-
-LogicalResult
-VulkanLaunchFuncToVulkanCallsPass::deduceMemRefRank(Value launchCallArg,
- uint32_t &rank) {
- // Deduce the rank from the type used to allocate the lowered MemRef.
- auto alloca = launchCallArg.getDefiningOp<LLVM::AllocaOp>();
- if (!alloca)
- return failure();
-
- std::optional<Type> elementType = alloca.getElemType();
- assert(elementType && "expected to work with opaque pointers");
- auto llvmDescriptorTy = dyn_cast<LLVM::LLVMStructType>(*elementType);
- // template <typename Elem, size_t Rank>
- // struct {
- // Elem *allocated;
- // Elem *aligned;
- // int64_t offset;
- // int64_t sizes[Rank]; // omitted when rank == 0
- // int64_t strides[Rank]; // omitted when rank == 0
- // };
- if (!llvmDescriptorTy)
- return failure();
-
- if (llvmDescriptorTy.getBody().size() == 3) {
- rank = 0;
- return success();
- }
- rank =
- cast<LLVM::LLVMArrayType>(llvmDescriptorTy.getBody()[3]).getNumElements();
- return success();
-}
-
-void VulkanLaunchFuncToVulkanCallsPass::declareVulkanFunctions(Location loc) {
- ModuleOp module = getOperation();
- auto builder = OpBuilder::atBlockEnd(module.getBody());
-
- if (!module.lookupSymbol(kSetEntryPoint)) {
- builder.create<LLVM::LLVMFuncOp>(
- loc, kSetEntryPoint,
- LLVM::LLVMFunctionType::get(getVoidType(),
- {getPointerType(), getPointerType()}));
- }
-
- if (!module.lookupSymbol(kSetNumWorkGroups)) {
- builder.create<LLVM::LLVMFuncOp>(
- loc, kSetNumWorkGroups,
- LLVM::LLVMFunctionType::get(getVoidType(),
- {getPointerType(), getInt64Type(),
- getInt64Type(), getInt64Type()}));
- }
-
- if (!module.lookupSymbol(kSetBinaryShader)) {
- builder.create<LLVM::LLVMFuncOp>(
- loc, kSetBinaryShader,
- LLVM::LLVMFunctionType::get(
- getVoidType(),
- {getPointerType(), getPointerType(), getInt32Type()}));
- }
-
- if (!module.lookupSymbol(kRunOnVulkan)) {
- builder.create<LLVM::LLVMFuncOp>(
- loc, kRunOnVulkan,
- LLVM::LLVMFunctionType::get(getVoidType(), {getPointerType()}));
- }
-
- for (unsigned i = 1; i <= 3; i++) {
- SmallVector<Type, 5> types{
- Float32Type::get(&getContext()), IntegerType::get(&getContext(), 32),
- IntegerType::get(&getContext(), 16), IntegerType::get(&getContext(), 8),
- Float16Type::get(&getContext())};
- for (auto type : types) {
- std::string fnName = "bindMemRef" + std::to_string(i) + "D" +
- std::string(stringifyType(type));
- if (isa<Float16Type>(type))
- type = IntegerType::get(&getContext(), 16);
- if (!module.lookupSymbol(fnName)) {
- auto fnType = LLVM::LLVMFunctionType::get(
- getVoidType(),
- {llvmPointerType, getInt32Type(), getInt32Type(), llvmPointerType},
- /*isVarArg=*/false);
- builder.create<LLVM::LLVMFuncOp>(loc, fnName, fnType);
- }
- }
- }
-
- if (!module.lookupSymbol(kInitVulkan)) {
- builder.create<LLVM::LLVMFuncOp>(
- loc, kInitVulkan, LLVM::LLVMFunctionType::get(getPointerType(), {}));
- }
-
- if (!module.lookupSymbol(kDeinitVulkan)) {
- builder.create<LLVM::LLVMFuncOp>(
- loc, kDeinitVulkan,
- LLVM::LLVMFunctionType::get(getVoidType(), {getPointerType()}));
- }
-}
-
-Value VulkanLaunchFuncToVulkanCallsPass::createEntryPointNameConstant(
- StringRef name, Location loc, OpBuilder &builder) {
- SmallString<16> shaderName(name.begin(), name.end());
- // Append `\0` to follow C style string given that LLVM::createGlobalString()
- // won't handle this directly for us.
- shaderName.push_back('\0');
-
- std::string entryPointGlobalName = (name + "_spv_entry_point_name").str();
- return LLVM::createGlobalString(loc, builder, entryPointGlobalName,
- shaderName, LLVM::Linkage::Internal);
-}
-
-void VulkanLaunchFuncToVulkanCallsPass::translateVulkanLaunchCall(
- LLVM::CallOp cInterfaceVulkanLaunchCallOp) {
- OpBuilder builder(cInterfaceVulkanLaunchCallOp);
- Location loc = cInterfaceVulkanLaunchCallOp.getLoc();
- // Create call to `initVulkan`.
- auto initVulkanCall = builder.create<LLVM::CallOp>(
- loc, TypeRange{getPointerType()}, kInitVulkan);
- // The result of `initVulkan` function is a pointer to Vulkan runtime, we
- // need to pass that pointer to each Vulkan runtime call.
- auto vulkanRuntime = initVulkanCall.getResult();
-
- // Create LLVM global with SPIR-V binary data, so we can pass a pointer with
- // that data to runtime call.
- Value ptrToSPIRVBinary = LLVM::createGlobalString(
- loc, builder, kSPIRVBinary, spirvAttributes.blob.getValue(),
- LLVM::Linkage::Internal);
-
- // Create LLVM constant for the size of SPIR-V binary shader.
- Value binarySize = builder.create<LLVM::ConstantOp>(
- loc, getInt32Type(), spirvAttributes.blob.getValue().size());
-
- // Create call to `bindMemRef` for each memref operand.
- createBindMemRefCalls(cInterfaceVulkanLaunchCallOp, vulkanRuntime);
-
- // Create call to `setBinaryShader` runtime function with the given pointer to
- // SPIR-V binary and binary size.
- builder.create<LLVM::CallOp>(
- loc, TypeRange(), kSetBinaryShader,
- ValueRange{vulkanRuntime, ptrToSPIRVBinary, binarySize});
- // Create LLVM global with entry point name.
- Value entryPointName = createEntryPointNameConstant(
- spirvAttributes.entryPoint.getValue(), loc, builder);
- // Create call to `setEntryPoint` runtime function with the given pointer to
- // entry point name.
- builder.create<LLVM::CallOp>(loc, TypeRange(), kSetEntryPoint,
- ValueRange{vulkanRuntime, entryPointName});
-
- // Create number of local workgroup for each dimension.
- builder.create<LLVM::CallOp>(
- loc, TypeRange(), kSetNumWorkGroups,
- ValueRange{vulkanRuntime, cInterfaceVulkanLaunchCallOp.getOperand(0),
- cInterfaceVulkanLaunchCallOp.getOperand(1),
- cInterfaceVulkanLaunchCallOp.getOperand(2)});
-
- // Create call to `runOnVulkan` runtime function.
- builder.create<LLVM::CallOp>(loc, TypeRange(), kRunOnVulkan,
- ValueRange{vulkanRuntime});
-
- // Create call to 'deinitVulkan' runtime function.
- builder.create<LLVM::CallOp>(loc, TypeRange(), kDeinitVulkan,
- ValueRange{vulkanRuntime});
-
- // Declare runtime functions.
- declareVulkanFunctions(loc);
-
- cInterfaceVulkanLaunchCallOp.erase();
-}
diff --git a/mlir/lib/ExecutionEngine/CMakeLists.txt b/mlir/lib/ExecutionEngine/CMakeLists.txt
index cf44a02cf5cb94..1579cc36b5d065 100644
--- a/mlir/lib/ExecutionEngine/CMakeLists.txt
+++ b/mlir/lib/ExecutionEngine/CMakeLists.txt
@@ -16,6 +16,9 @@ set(LLVM_OPTIONAL_SOURCES
JitRunner.cpp
SpirvCpuRuntimeWrappers.cpp
SyclRuntimeWrappers.cpp
+ VulkanRuntimeWrappers.cpp
+ VulkanRuntime.cpp
+ VulkanRuntime.h
)
# Use a separate library for OptUtils, to avoid pulling in the entire JIT and
@@ -415,4 +418,46 @@ if(LLVM_ENABLE_PIC)
PRIVATE
mlir_spirv_cpu_runtime_EXPORTS)
endif()
+
+ if (MLIR_ENABLE_VULKAN_RUNNER)
+ find_package(Vulkan)
+
+ # If Vulkan is not found try a path specified by VULKAN_SDK.
+ if (NOT Vulkan_FOUND)
+ if ("$ENV{VULKAN_SDK}" STREQUAL "")
+ message(FATAL_ERROR "Vulkan not found through CMake; please provide "
+ "VULKAN_SDK path as an environment variable")
+ endif()
+
+ find_library(Vulkan_LIBRARY vulkan HINTS "$ENV{VULKAN_SDK}/lib" REQUIRED)
+ if (Vulkan_LIBRARY)
+ set(Vulkan_FOUND ON)
+ set(Vulkan_INCLUDE_DIR "$ENV{VULKAN_SDK}/include")
+ message(STATUS "Found Vulkan: " ${Vulkan_LIBRARY})
+ endif()
+ endif()
+
+ if (NOT Vulkan_FOUND)
+ message(FATAL_ERROR "Cannot find Vulkan library")
+ endif()
+
+ add_llvm_library(mlir_vulkan_runtime SHARED
+ VulkanRuntimeWrappers.cpp
+ VulkanRuntime.cpp
+ )
+
+ target_include_directories(mlir_vulkan_runtime
+ PUBLIC
+ ${Vulkan_INCLUDE_DIR}
+ )
+
+ # *IMPORTANT*: This library cannot depend on LLVM libraries. Otherwise,
+ # it may cause LLVM version conflict when used together with other shared
+ # libraries depending on LLVM. Notably, Mesa, who implements Vulkan
+ # drivers on Linux, depends on the system libLLVM.so.
+ target_link_libraries(mlir_vulkan_runtime
+ PUBLIC
+ ${Vulkan_LIBRARY}
+ )
+ endif()
endif()
diff --git a/mlir/tools/mlir-vulkan-runner/VulkanRuntime.cpp b/mlir/lib/ExecutionEngine/VulkanRuntime.cpp
similarity index 100%
rename from mlir/tools/mlir-vulkan-runner/VulkanRuntime.cpp
rename to mlir/lib/ExecutionEngine/VulkanRuntime.cpp
diff --git a/mlir/tools/mlir-vulkan-runner/VulkanRuntime.h b/mlir/lib/ExecutionEngine/VulkanRuntime.h
similarity index 100%
rename from mlir/tools/mlir-vulkan-runner/VulkanRuntime.h
rename to mlir/lib/ExecutionEngine/VulkanRuntime.h
diff --git a/mlir/tools/mlir-vulkan-runner/vulkan-runtime-wrappers.cpp b/mlir/lib/ExecutionEngine/VulkanRuntimeWrappers.cpp
similarity index 75%
rename from mlir/tools/mlir-vulkan-runner/vulkan-runtime-wrappers.cpp
rename to mlir/lib/ExecutionEngine/VulkanRuntimeWrappers.cpp
index 8d1bac3b6f2869..c414fe53b2f3c2 100644
--- a/mlir/tools/mlir-vulkan-runner/vulkan-runtime-wrappers.cpp
+++ b/mlir/lib/ExecutionEngine/VulkanRuntimeWrappers.cpp
@@ -1,4 +1,4 @@
-//===- vulkan-runtime-wrappers.cpp - MLIR Vulkan runner wrapper library ---===//
+//===- VulkanRuntimeWrappers.cpp - MLIR Vulkan runner wrapper library -----===//
//
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
// See https://llvm.org/LICENSE.txt for license information.
@@ -113,23 +113,12 @@ struct MemRefDescriptor {
int64_t strides[N];
};
-template <typename T, uint32_t S>
-void bindMemRef(void *vkRuntimeManager, DescriptorSetIndex setIndex,
- BindingIndex bindIndex, MemRefDescriptor<T, S> *ptr) {
- uint32_t size = sizeof(T);
- for (unsigned i = 0; i < S; i++)
- size *= ptr->sizes[i];
- VulkanHostMemoryBuffer memBuffer{ptr->aligned, size};
- reinterpret_cast<VulkanRuntimeManager *>(vkRuntimeManager)
- ->setResourceData(setIndex, bindIndex, memBuffer);
-}
-
extern "C" {
//===----------------------------------------------------------------------===//
//
-// New wrappers, intended for mlir-cpu-runner. Calls to these are generated by
-// GPUToLLVMConversionPass.
+// Wrappers intended for mlir-cpu-runner. Uses of GPU dialect operations get
+// lowered to calls to these functions by GPUToLLVMConversionPass.
//
//===----------------------------------------------------------------------===//
@@ -203,68 +192,10 @@ mgpuLaunchKernel(void *vkKernel, size_t gridX, size_t gridY, size_t gridZ,
//===----------------------------------------------------------------------===//
//
-// Old wrappers, intended for mlir-vulkan-runner. Calls to these are generated
-// by LaunchFuncToVulkanCallsPass.
+// Miscellaneous utility functions that can be directly used by tests.
//
//===----------------------------------------------------------------------===//
-/// Initializes `VulkanRuntimeManager` and returns a pointer to it.
-VULKAN_WRAPPER_SYMBOL_EXPORT void *initVulkan() {
- return new VulkanRuntimeManager();
-}
-
-/// Deinitializes `VulkanRuntimeManager` by the given pointer.
-VULKAN_WRAPPER_SYMBOL_EXPORT void deinitVulkan(void *vkRuntimeManager) {
- delete reinterpret_cast<VulkanRuntimeManager *>(vkRuntimeManager);
-}
-
-VULKAN_WRAPPER_SYMBOL_EXPORT void runOnVulkan(void *vkRuntimeManager) {
- reinterpret_cast<VulkanRuntimeManager *>(vkRuntimeManager)->runOnVulkan();
-}
-
-VULKAN_WRAPPER_SYMBOL_EXPORT void setEntryPoint(void *vkRuntimeManager,
- const char *entryPoint) {
- reinterpret_cast<VulkanRuntimeManager *>(vkRuntimeManager)
- ->setEntryPoint(entryPoint);
-}
-
-VULKAN_WRAPPER_SYMBOL_EXPORT void
-setNumWorkGroups(void *vkRuntimeManager, uint32_t x, uint32_t y, uint32_t z) {
- reinterpret_cast<VulkanRuntimeManager *>(vkRuntimeManager)
- ->setNumWorkGroups({x, y, z});
-}
-
-VULKAN_WRAPPER_SYMBOL_EXPORT void
-setBinaryShader(void *vkRuntimeManager, uint8_t *shader, uint32_t size) {
- reinterpret_cast<VulkanRuntimeManager *>(vkRuntimeManager)
- ->setShaderModule(shader, size);
-}
-
-/// Binds the given memref to the given descriptor set and descriptor
-/// index.
-#define DECLARE_BIND_MEMREF(size, type, typeName) \
- VULKAN_WRAPPER_SYMBOL_EXPORT void bindMemRef##size##D##typeName( \
- void *vkRuntimeManager, DescriptorSetIndex setIndex, \
- BindingIndex bindIndex, MemRefDescriptor<type, size> *ptr) { \
- bindMemRef<type, size>(vkRuntimeManager, setIndex, bindIndex, ptr); \
- }
-
-DECLARE_BIND_MEMREF(1, float, Float)
-DECLARE_BIND_MEMREF(2, float, Float)
-DECLARE_BIND_MEMREF(3, float, Float)
-DECLARE_BIND_MEMREF(1, int32_t, Int32)
-DECLARE_BIND_MEMREF(2, int32_t, Int32)
-DECLARE_BIND_MEMREF(3, int32_t, Int32)
-DECLARE_BIND_MEMREF(1, int16_t, Int16)
-DECLARE_BIND_MEMREF(2, int16_t, Int16)
-DECLARE_BIND_MEMREF(3, int16_t, Int16)
-DECLARE_BIND_MEMREF(1, int8_t, Int8)
-DECLARE_BIND_MEMREF(2, int8_t, Int8)
-DECLARE_BIND_MEMREF(3, int8_t, Int8)
-DECLARE_BIND_MEMREF(1, int16_t, Half)
-DECLARE_BIND_MEMREF(2, int16_t, Half)
-DECLARE_BIND_MEMREF(3, int16_t, Half)
-
/// Fills the given 1D float memref with the given float value.
VULKAN_WRAPPER_SYMBOL_EXPORT void
_mlir_ciface_fillResource1DFloat(MemRefDescriptor<float, 1> *ptr, // NOLINT
diff --git a/mlir/test/CMakeLists.txt b/mlir/test/CMakeLists.txt
index 58d16a657297e6..69c2e597868923 100644
--- a/mlir/test/CMakeLists.txt
+++ b/mlir/test/CMakeLists.txt
@@ -206,7 +206,7 @@ endif()
if(MLIR_ENABLE_VULKAN_RUNNER)
list(APPEND MLIR_TEST_DEPENDS
- mlir-vulkan-runner
+ mlir_vulkan_runtime
)
endif()
diff --git a/mlir/test/Conversion/GPUToVulkan/invoke-vulkan.mlir b/mlir/test/Conversion/GPUToVulkan/invoke-vulkan.mlir
deleted file mode 100644
index fe8a36ee29a9f0..00000000000000
--- a/mlir/test/Conversion/GPUToVulkan/invoke-vulkan.mlir
+++ /dev/null
@@ -1,62 +0,0 @@
-// RUN: mlir-opt %s -launch-func-to-vulkan | FileCheck %s
-
-// CHECK: llvm.mlir.global internal constant @kernel_spv_entry_point_name
-// CHECK: llvm.mlir.global internal constant @SPIRV_BIN
-// CHECK: %[[Vulkan_Runtime_ptr:.*]] = llvm.call @initVulkan() : () -> !llvm.ptr
-// CHECK: %[[addressof_SPIRV_BIN:.*]] = llvm.mlir.addressof @SPIRV_BIN
-// CHECK: %[[SPIRV_BIN_ptr:.*]] = llvm.getelementptr %[[addressof_SPIRV_BIN]]
-// CHECK: %[[SPIRV_BIN_size:.*]] = llvm.mlir.constant
-// CHECK: llvm.call @bindMemRef1DFloat(%[[Vulkan_Runtime_ptr]], %{{.*}}, %{{.*}}, %{{.*}}) : (!llvm.ptr, i32, i32, !llvm.ptr) -> ()
-// CHECK: llvm.call @setBinaryShader(%[[Vulkan_Runtime_ptr]], %[[SPIRV_BIN_ptr]], %[[SPIRV_BIN_size]]) : (!llvm.ptr, !llvm.ptr, i32) -> ()
-// CHECK: %[[addressof_entry_point:.*]] = llvm.mlir.addressof @kernel_spv_entry_point_name
-// CHECK: %[[entry_point_ptr:.*]] = llvm.getelementptr %[[addressof_entry_point]]
-// CHECK: llvm.call @setEntryPoint(%[[Vulkan_Runtime_ptr]], %[[entry_point_ptr]]) : (!llvm.ptr, !llvm.ptr) -> ()
-// CHECK: llvm.call @setNumWorkGroups(%[[Vulkan_Runtime_ptr]], %{{.*}}, %{{.*}}, %{{.*}}) : (!llvm.ptr, i64, i64, i64) -> ()
-// CHECK: llvm.call @runOnVulkan(%[[Vulkan_Runtime_ptr]]) : (!llvm.ptr) -> ()
-// CHECK: llvm.call @deinitVulkan(%[[Vulkan_Runtime_ptr]]) : (!llvm.ptr) -> ()
-
-// CHECK: llvm.func @bindMemRef1DHalf(!llvm.ptr, i32, i32, !llvm.ptr)
-
-module attributes {gpu.container_module} {
- llvm.func @malloc(i64) -> !llvm.ptr
- llvm.func @foo() {
- %0 = llvm.mlir.constant(12 : index) : i64
- %1 = llvm.mlir.zero : !llvm.ptr
- %2 = llvm.mlir.constant(1 : index) : i64
- %3 = llvm.getelementptr %1[%2] : (!llvm.ptr, i64) -> !llvm.ptr, f32
- %4 = llvm.ptrtoint %3 : !llvm.ptr to i64
- %5 = llvm.mul %0, %4 : i64
- %6 = llvm.call @malloc(%5) : (i64) -> !llvm.ptr
- %8 = llvm.mlir.undef : !llvm.struct<(ptr, ptr, i64, array<1 x i64>, array<1 x i64>)>
- %9 = llvm.insertvalue %6, %8[0] : !llvm.struct<(ptr, ptr, i64, array<1 x i64>, array<1 x i64>)>
- %10 = llvm.insertvalue %6, %9[1] : !llvm.struct<(ptr, ptr, i64, array<1 x i64>, array<1 x i64>)>
- %11 = llvm.mlir.constant(0 : index) : i64
- %12 = llvm.insertvalue %11, %10[2] : !llvm.struct<(ptr, ptr, i64, array<1 x i64>, array<1 x i64>)>
- %13 = llvm.mlir.constant(1 : index) : i64
- %14 = llvm.insertvalue %0, %12[3, 0] : !llvm.struct<(ptr, ptr, i64, array<1 x i64>, array<1 x i64>)>
- %15 = llvm.insertvalue %13, %14[4, 0] : !llvm.struct<(ptr, ptr, i64, array<1 x i64>, array<1 x i64>)>
- %16 = llvm.mlir.constant(1 : index) : i64
- %17 = llvm.extractvalue %15[0] : !llvm.struct<(ptr, ptr, i64, array<1 x i64>, array<1 x i64>)>
- %18 = llvm.extractvalue %15[1] : !llvm.struct<(ptr, ptr, i64, array<1 x i64>, array<1 x i64>)>
- %19 = llvm.extractvalue %15[2] : !llvm.struct<(ptr, ptr, i64, array<1 x i64>, array<1 x i64>)>
- %20 = llvm.extractvalue %15[3, 0] : !llvm.struct<(ptr, ptr, i64, array<1 x i64>, array<1 x i64>)>
- %21 = llvm.extractvalue %15[4, 0] : !llvm.struct<(ptr, ptr, i64, array<1 x i64>, array<1 x i64>)>
- llvm.call @vulkanLaunch(%16, %16, %16, %17, %18, %19, %20, %21) {spirv_blob = "\03\02#\07\00", spirv_element_types = [f32], spirv_entry_point = "kernel"}
- : (i64, i64, i64, !llvm.ptr, !llvm.ptr, i64, i64, i64) -> ()
- llvm.return
- }
- llvm.func @vulkanLaunch(%arg0: i64, %arg1: i64, %arg2: i64, %arg6: !llvm.ptr, %arg7: !llvm.ptr, %arg8: i64, %arg9: i64, %arg10: i64) {
- %0 = llvm.mlir.undef : !llvm.struct<(ptr, ptr, i64, array<1 x i64>, array<1 x i64>)>
- %1 = llvm.insertvalue %arg6, %0[0] : !llvm.struct<(ptr, ptr, i64, array<1 x i64>, array<1 x i64>)>
- %2 = llvm.insertvalue %arg7, %1[1] : !llvm.struct<(ptr, ptr, i64, array<1 x i64>, array<1 x i64>)>
- %3 = llvm.insertvalue %arg8, %2[2] : !llvm.struct<(ptr, ptr, i64, array<1 x i64>, array<1 x i64>)>
- %4 = llvm.insertvalue %arg9, %3[3, 0] : !llvm.struct<(ptr, ptr, i64, array<1 x i64>, array<1 x i64>)>
- %5 = llvm.insertvalue %arg10, %4[4, 0] : !llvm.struct<(ptr, ptr, i64, array<1 x i64>, array<1 x i64>)>
- %6 = llvm.mlir.constant(1 : index) : i64
- %7 = llvm.alloca %6 x !llvm.struct<(ptr, ptr, i64, array<1 x i64>, array<1 x i64>)> : (i64) -> !llvm.ptr
- llvm.store %5, %7 : !llvm.struct<(ptr, ptr, i64, array<1 x i64>, array<1 x i64>)>, !llvm.ptr
- llvm.call @_mlir_ciface_vulkanLaunch(%arg0, %arg1, %arg2, %7) : (i64, i64, i64, !llvm.ptr) -> ()
- llvm.return
- }
- llvm.func @_mlir_ciface_vulkanLaunch(i64, i64, i64, !llvm.ptr)
-}
diff --git a/mlir/test/Conversion/GPUToVulkan/lower-gpu-launch-vulkan-launch.mlir b/mlir/test/Conversion/GPUToVulkan/lower-gpu-launch-vulkan-launch.mlir
deleted file mode 100644
index 96ee1866517e6d..00000000000000
--- a/mlir/test/Conversion/GPUToVulkan/lower-gpu-launch-vulkan-launch.mlir
+++ /dev/null
@@ -1,35 +0,0 @@
-// 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} {
- 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
- }
- }
- func.func @foo() {
- %0 = memref.alloc() : memref<12xf32>
- %c1 = arith.constant 1 : index
- gpu.launch_func @kernels::@kernel
- blocks in(%c1, %c1, %c1)
- threads in(%c1, %c1, %c1)
- args(%0 : memref<12xf32>)
- return
- }
-}
diff --git a/mlir/test/mlir-vulkan-runner/addf.mlir b/mlir/test/Integration/GPU/Vulkan/addf.mlir
similarity index 91%
rename from mlir/test/mlir-vulkan-runner/addf.mlir
rename to mlir/test/Integration/GPU/Vulkan/addf.mlir
index 71f87a8b0d5c82..f4d2463d413b84 100644
--- a/mlir/test/mlir-vulkan-runner/addf.mlir
+++ b/mlir/test/Integration/GPU/Vulkan/addf.mlir
@@ -1,5 +1,5 @@
-// RUN: mlir-opt %s -test-vulkan-runner-pipeline=to-llvm \
-// RUN: | mlir-cpu-runner - --shared-libs=%vulkan-runtime-wrappers,%mlir_runner_utils --entry-point-result=void | FileCheck %s
+// RUN: mlir-opt %s -test-vulkan-runner-pipeline \
+// RUN: | mlir-cpu-runner - --shared-libs=%mlir_vulkan_runtime,%mlir_runner_utils --entry-point-result=void | FileCheck %s
// CHECK: [3.3, 3.3, 3.3, 3.3, 3.3, 3.3, 3.3, 3.3]
module attributes {
diff --git a/mlir/test/mlir-vulkan-runner/addf_if.mlir b/mlir/test/Integration/GPU/Vulkan/addf_if.mlir
similarity index 91%
rename from mlir/test/mlir-vulkan-runner/addf_if.mlir
rename to mlir/test/Integration/GPU/Vulkan/addf_if.mlir
index 6fe51a83482dca..2512d0f8e6efa4 100644
--- a/mlir/test/mlir-vulkan-runner/addf_if.mlir
+++ b/mlir/test/Integration/GPU/Vulkan/addf_if.mlir
@@ -1,5 +1,5 @@
-// RUN: mlir-opt %s -test-vulkan-runner-pipeline=to-llvm \
-// RUN: | mlir-cpu-runner - --shared-libs=%vulkan-runtime-wrappers,%mlir_runner_utils --entry-point-result=void | FileCheck %s
+// RUN: mlir-opt %s -test-vulkan-runner-pipeline \
+// RUN: | mlir-cpu-runner - --shared-libs=%mlir_vulkan_runtime,%mlir_runner_utils --entry-point-result=void | FileCheck %s
// CHECK: [3.3, 3.3, 3.3, 3.3, 0, 0, 0, 0]
module attributes {
diff --git a/mlir/test/mlir-vulkan-runner/addi.mlir b/mlir/test/Integration/GPU/Vulkan/addi.mlir
similarity index 91%
rename from mlir/test/mlir-vulkan-runner/addi.mlir
rename to mlir/test/Integration/GPU/Vulkan/addi.mlir
index e60eb7696770a8..abf695c61f3b30 100644
--- a/mlir/test/mlir-vulkan-runner/addi.mlir
+++ b/mlir/test/Integration/GPU/Vulkan/addi.mlir
@@ -1,5 +1,5 @@
-// RUN: mlir-opt %s -test-vulkan-runner-pipeline=to-llvm \
-// RUN: | mlir-cpu-runner - --shared-libs=%vulkan-runtime-wrappers,%mlir_runner_utils --entry-point-result=void | FileCheck %s
+// RUN: mlir-opt %s -test-vulkan-runner-pipeline \
+// RUN: | mlir-cpu-runner - --shared-libs=%mlir_vulkan_runtime,%mlir_runner_utils --entry-point-result=void | FileCheck %s
// CHECK-COUNT-64: [3, 3, 3, 3, 3, 3, 3, 3]
module attributes {
diff --git a/mlir/test/mlir-vulkan-runner/addi8.mlir b/mlir/test/Integration/GPU/Vulkan/addi8.mlir
similarity index 92%
rename from mlir/test/mlir-vulkan-runner/addi8.mlir
rename to mlir/test/Integration/GPU/Vulkan/addi8.mlir
index 6374f29fd0ed2d..fd43422fbafad4 100644
--- a/mlir/test/mlir-vulkan-runner/addi8.mlir
+++ b/mlir/test/Integration/GPU/Vulkan/addi8.mlir
@@ -1,5 +1,5 @@
-// RUN: mlir-opt %s -test-vulkan-runner-pipeline=to-llvm \
-// RUN: | mlir-cpu-runner - --shared-libs=%vulkan-runtime-wrappers,%mlir_runner_utils --entry-point-result=void | FileCheck %s
+// RUN: mlir-opt %s -test-vulkan-runner-pipeline \
+// RUN: | mlir-cpu-runner - --shared-libs=%mlir_vulkan_runtime,%mlir_runner_utils --entry-point-result=void | FileCheck %s
// CHECK-COUNT-64: [3, 3, 3, 3, 3, 3, 3, 3]
module attributes {
diff --git a/mlir/test/mlir-vulkan-runner/addui_extended.mlir b/mlir/test/Integration/GPU/Vulkan/addui_extended.mlir
similarity index 90%
rename from mlir/test/mlir-vulkan-runner/addui_extended.mlir
rename to mlir/test/Integration/GPU/Vulkan/addui_extended.mlir
index 0894bc301f2e3c..d048b7b0290a9e 100644
--- a/mlir/test/mlir-vulkan-runner/addui_extended.mlir
+++ b/mlir/test/Integration/GPU/Vulkan/addui_extended.mlir
@@ -1,14 +1,14 @@
// Make sure that addition with carry produces expected results
// with and without expansion to primitive add/cmp ops for WebGPU.
-// RUN: mlir-opt %s -test-vulkan-runner-pipeline=to-llvm \
+// RUN: mlir-opt %s -test-vulkan-runner-pipeline \
// RUN: | mlir-cpu-runner - \
-// RUN: --shared-libs=%vulkan-runtime-wrappers,%mlir_runner_utils \
+// RUN: --shared-libs=%mlir_vulkan_runtime,%mlir_runner_utils \
// RUN: --entry-point-result=void | FileCheck %s
-// RUN: mlir-opt %s -test-vulkan-runner-pipeline="spirv-webgpu-prepare to-llvm" \
+// RUN: mlir-opt %s -test-vulkan-runner-pipeline=spirv-webgpu-prepare \
// RUN: | mlir-cpu-runner - \
-// RUN: --shared-libs=%vulkan-runtime-wrappers,%mlir_runner_utils \
+// RUN: --shared-libs=%mlir_vulkan_runtime,%mlir_runner_utils \
// RUN: --entry-point-result=void | FileCheck %s
// CHECK: [0, 42, 0, 42]
diff --git a/mlir/test/mlir-vulkan-runner/lit.local.cfg b/mlir/test/Integration/GPU/Vulkan/lit.local.cfg
similarity index 100%
rename from mlir/test/mlir-vulkan-runner/lit.local.cfg
rename to mlir/test/Integration/GPU/Vulkan/lit.local.cfg
diff --git a/mlir/test/mlir-vulkan-runner/mulf.mlir b/mlir/test/Integration/GPU/Vulkan/mulf.mlir
similarity index 91%
rename from mlir/test/mlir-vulkan-runner/mulf.mlir
rename to mlir/test/Integration/GPU/Vulkan/mulf.mlir
index bd1b8d9abf4c9f..f1f71bca73c3e6 100644
--- a/mlir/test/mlir-vulkan-runner/mulf.mlir
+++ b/mlir/test/Integration/GPU/Vulkan/mulf.mlir
@@ -1,5 +1,5 @@
-// RUN: mlir-opt %s -test-vulkan-runner-pipeline=to-llvm \
-// RUN: | mlir-cpu-runner - --shared-libs=%vulkan-runtime-wrappers,%mlir_runner_utils --entry-point-result=void | FileCheck %s
+// RUN: mlir-opt %s -test-vulkan-runner-pipeline \
+// RUN: | mlir-cpu-runner - --shared-libs=%mlir_vulkan_runtime,%mlir_runner_utils --entry-point-result=void | FileCheck %s
// CHECK-COUNT-4: [6, 6, 6, 6]
module attributes {
diff --git a/mlir/test/mlir-vulkan-runner/smul_extended.mlir b/mlir/test/Integration/GPU/Vulkan/smul_extended.mlir
similarity index 90%
rename from mlir/test/mlir-vulkan-runner/smul_extended.mlir
rename to mlir/test/Integration/GPU/Vulkan/smul_extended.mlir
index 0ef86f46562e86..ac46b9035c13c6 100644
--- a/mlir/test/mlir-vulkan-runner/smul_extended.mlir
+++ b/mlir/test/Integration/GPU/Vulkan/smul_extended.mlir
@@ -1,14 +1,14 @@
// Make sure that signed extended multiplication produces expected results
// with and without expansion to primitive mul/add ops for WebGPU.
-// RUN: mlir-opt %s -test-vulkan-runner-pipeline=to-llvm \
+// RUN: mlir-opt %s -test-vulkan-runner-pipeline \
// RUN: | mlir-cpu-runner - \
-// RUN: --shared-libs=%vulkan-runtime-wrappers,%mlir_runner_utils \
+// RUN: --shared-libs=%mlir_vulkan_runtime,%mlir_runner_utils \
// RUN: --entry-point-result=void | FileCheck %s
-// RUN: mlir-opt %s -test-vulkan-runner-pipeline="spirv-webgpu-prepare to-llvm" \
+// RUN: mlir-opt %s -test-vulkan-runner-pipeline=spirv-webgpu-prepare \
// RUN: | mlir-cpu-runner - \
-// RUN: --shared-libs=%vulkan-runtime-wrappers,%mlir_runner_utils \
+// RUN: --shared-libs=%mlir_vulkan_runtime,%mlir_runner_utils \
// RUN: --entry-point-result=void | FileCheck %s
// CHECK: [0, 1, -2, 1, 1048560, -87620295, -131071, 560969770]
diff --git a/mlir/test/mlir-vulkan-runner/subf.mlir b/mlir/test/Integration/GPU/Vulkan/subf.mlir
similarity index 92%
rename from mlir/test/mlir-vulkan-runner/subf.mlir
rename to mlir/test/Integration/GPU/Vulkan/subf.mlir
index e8cee0e021a277..50c63abc8c10a4 100644
--- a/mlir/test/mlir-vulkan-runner/subf.mlir
+++ b/mlir/test/Integration/GPU/Vulkan/subf.mlir
@@ -1,5 +1,5 @@
-// RUN: mlir-opt %s -test-vulkan-runner-pipeline=to-llvm \
-// RUN: | mlir-cpu-runner - --shared-libs=%vulkan-runtime-wrappers,%mlir_runner_utils --entry-point-result=void | FileCheck %s
+// RUN: mlir-opt %s -test-vulkan-runner-pipeline \
+// RUN: | mlir-cpu-runner - --shared-libs=%mlir_vulkan_runtime,%mlir_runner_utils --entry-point-result=void | FileCheck %s
// CHECK-COUNT-32: [2.2, 2.2, 2.2, 2.2]
module attributes {
diff --git a/mlir/test/mlir-vulkan-runner/time.mlir b/mlir/test/Integration/GPU/Vulkan/time.mlir
similarity index 92%
rename from mlir/test/mlir-vulkan-runner/time.mlir
rename to mlir/test/Integration/GPU/Vulkan/time.mlir
index f6284478742389..f506f6be15df21 100644
--- a/mlir/test/mlir-vulkan-runner/time.mlir
+++ b/mlir/test/Integration/GPU/Vulkan/time.mlir
@@ -1,5 +1,5 @@
-// RUN: mlir-opt %s -test-vulkan-runner-pipeline=to-llvm \
-// RUN: | mlir-cpu-runner - --shared-libs=%vulkan-runtime-wrappers,%mlir_runner_utils --entry-point-result=void | FileCheck %s
+// RUN: mlir-opt %s -test-vulkan-runner-pipeline \
+// RUN: | mlir-cpu-runner - --shared-libs=%mlir_vulkan_runtime,%mlir_runner_utils --entry-point-result=void | FileCheck %s
// CHECK: Compute shader execution time
// CHECK: Command buffer submit time
diff --git a/mlir/test/mlir-vulkan-runner/umul_extended.mlir b/mlir/test/Integration/GPU/Vulkan/umul_extended.mlir
similarity index 90%
rename from mlir/test/mlir-vulkan-runner/umul_extended.mlir
rename to mlir/test/Integration/GPU/Vulkan/umul_extended.mlir
index 5936c808435c19..0f01a884296931 100644
--- a/mlir/test/mlir-vulkan-runner/umul_extended.mlir
+++ b/mlir/test/Integration/GPU/Vulkan/umul_extended.mlir
@@ -1,14 +1,14 @@
// Make sure that unsigned extended multiplication produces expected results
// with and without expansion to primitive mul/add ops for WebGPU.
-// RUN: mlir-opt %s -test-vulkan-runner-pipeline=to-llvm \
+// RUN: mlir-opt %s -test-vulkan-runner-pipeline \
// RUN: | mlir-cpu-runner - \
-// RUN: --shared-libs=%vulkan-runtime-wrappers,%mlir_runner_utils \
+// RUN: --shared-libs=%mlir_vulkan_runtime,%mlir_runner_utils \
// RUN: --entry-point-result=void | FileCheck %s
-// RUN: mlir-opt %s -test-vulkan-runner-pipeline="spirv-webgpu-prepare to-llvm" \
+// RUN: mlir-opt %s -test-vulkan-runner-pipeline=spirv-webgpu-prepare \
// RUN: | mlir-cpu-runner - \
-// RUN: --shared-libs=%vulkan-runtime-wrappers,%mlir_runner_utils \
+// RUN: --shared-libs=%mlir_vulkan_runtime,%mlir_runner_utils \
// RUN: --entry-point-result=void | FileCheck %s
// CHECK: [0, 1, -2, 1, 1048560, -87620295, -131071, -49]
diff --git a/mlir/test/mlir-vulkan-runner/vector-deinterleave.mlir b/mlir/test/Integration/GPU/Vulkan/vector-deinterleave.mlir
similarity index 96%
rename from mlir/test/mlir-vulkan-runner/vector-deinterleave.mlir
rename to mlir/test/Integration/GPU/Vulkan/vector-deinterleave.mlir
index ebeb19cd6bcc5f..4e3f6ace15f6b2 100644
--- a/mlir/test/mlir-vulkan-runner/vector-deinterleave.mlir
+++ b/mlir/test/Integration/GPU/Vulkan/vector-deinterleave.mlir
@@ -1,6 +1,6 @@
-// RUN: mlir-opt %s -test-vulkan-runner-pipeline=to-llvm \
+// RUN: mlir-opt %s -test-vulkan-runner-pipeline \
// RUN: | mlir-cpu-runner - \
-// RUN: --shared-libs=%vulkan-runtime-wrappers,%mlir_runner_utils \
+// RUN: --shared-libs=%mlir_vulkan_runtime,%mlir_runner_utils \
// RUN: --entry-point-result=void | FileCheck %s
// CHECK: [0, 2]
diff --git a/mlir/test/mlir-vulkan-runner/vector-interleave.mlir b/mlir/test/Integration/GPU/Vulkan/vector-interleave.mlir
similarity index 96%
rename from mlir/test/mlir-vulkan-runner/vector-interleave.mlir
rename to mlir/test/Integration/GPU/Vulkan/vector-interleave.mlir
index 9314baf9b39c7e..f7f620bf766b83 100644
--- a/mlir/test/mlir-vulkan-runner/vector-interleave.mlir
+++ b/mlir/test/Integration/GPU/Vulkan/vector-interleave.mlir
@@ -1,6 +1,6 @@
-// RUN: mlir-opt %s -test-vulkan-runner-pipeline=to-llvm \
+// RUN: mlir-opt %s -test-vulkan-runner-pipeline \
// RUN: | mlir-cpu-runner - \
-// RUN: --shared-libs=%vulkan-runtime-wrappers,%mlir_runner_utils \
+// RUN: --shared-libs=%mlir_vulkan_runtime,%mlir_runner_utils \
// RUN: --entry-point-result=void | FileCheck %s
// CHECK: [0, 2, 1, 3]
diff --git a/mlir/test/mlir-vulkan-runner/vector-shuffle.mlir b/mlir/test/Integration/GPU/Vulkan/vector-shuffle.mlir
similarity index 96%
rename from mlir/test/mlir-vulkan-runner/vector-shuffle.mlir
rename to mlir/test/Integration/GPU/Vulkan/vector-shuffle.mlir
index cf3e2c569426b7..0f9c883091b89f 100644
--- a/mlir/test/mlir-vulkan-runner/vector-shuffle.mlir
+++ b/mlir/test/Integration/GPU/Vulkan/vector-shuffle.mlir
@@ -1,6 +1,6 @@
-// RUN: mlir-opt %s -test-vulkan-runner-pipeline=to-llvm \
+// RUN: mlir-opt %s -test-vulkan-runner-pipeline \
// RUN: | mlir-cpu-runner - \
-// RUN: --shared-libs=%vulkan-runtime-wrappers,%mlir_runner_utils \
+// RUN: --shared-libs=%mlir_vulkan_runtime,%mlir_runner_utils \
// RUN: --entry-point-result=void | FileCheck %s
// CHECK: [2, 1, 3, 3]
diff --git a/mlir/test/lib/Pass/TestVulkanRunnerPipeline.cpp b/mlir/test/lib/Pass/TestVulkanRunnerPipeline.cpp
index fa6a5a7140d8ca..e4cbbeb1f99bc4 100644
--- a/mlir/test/lib/Pass/TestVulkanRunnerPipeline.cpp
+++ b/mlir/test/lib/Pass/TestVulkanRunnerPipeline.cpp
@@ -6,7 +6,7 @@
//
//===----------------------------------------------------------------------===//
//
-// Implements a pipeline for use by mlir-vulkan-runner tests.
+// Implements a pipeline for use by Vulkan runner tests.
//
//===----------------------------------------------------------------------===//
@@ -33,9 +33,6 @@ struct VulkanRunnerPipelineOptions
Option<bool> spirvWebGPUPrepare{
*this, "spirv-webgpu-prepare",
llvm::cl::desc("Run MLIR transforms used when targetting WebGPU")};
- Option<bool> toLlvm{*this, "to-llvm",
- llvm::cl::desc("Run MLIR transforms to lower host code "
- "to LLVM, intended for mlir-cpu-runner")};
};
void buildTestVulkanRunnerPipeline(OpPassManager &passManager,
@@ -64,17 +61,14 @@ void buildTestVulkanRunnerPipeline(OpPassManager &passManager,
passManager.addPass(createGpuModuleToBinaryPass());
- if (options.toLlvm) {
- passManager.addPass(createFinalizeMemRefToLLVMConversionPass());
- passManager.nest<func::FuncOp>().addPass(
- LLVM::createRequestCWrappersPass());
- // vulkan-runtime-wrappers.cpp requires these calling convention options.
- GpuToLLVMConversionPassOptions opt;
- opt.hostBarePtrCallConv = false;
- opt.kernelBarePtrCallConv = true;
- opt.kernelIntersperseSizeCallConv = true;
- passManager.addPass(createGpuToLLVMConversionPass(opt));
- }
+ passManager.addPass(createFinalizeMemRefToLLVMConversionPass());
+ passManager.nest<func::FuncOp>().addPass(LLVM::createRequestCWrappersPass());
+ // VulkanRuntimeWrappers.cpp requires these calling convention options.
+ GpuToLLVMConversionPassOptions opt;
+ opt.hostBarePtrCallConv = false;
+ opt.kernelBarePtrCallConv = true;
+ opt.kernelIntersperseSizeCallConv = true;
+ passManager.addPass(createGpuToLLVMConversionPass(opt));
}
} // namespace
@@ -83,8 +77,9 @@ namespace mlir::test {
void registerTestVulkanRunnerPipeline() {
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 or mlir-cpu-runner.",
+ "Runs a series of passes intended for Vulkan runner tests. Lowers GPU "
+ "dialect to LLVM dialect for the host and to serialized Vulkan SPIR-V "
+ "for the device.",
buildTestVulkanRunnerPipeline);
}
} // namespace mlir::test
diff --git a/mlir/test/lit.cfg.py b/mlir/test/lit.cfg.py
index c28623123d9991..14ffb97f0f065d 100644
--- a/mlir/test/lit.cfg.py
+++ b/mlir/test/lit.cfg.py
@@ -197,7 +197,7 @@ def find_real_python_interpreter():
]
if config.enable_vulkan_runner:
- tools.extend([add_runtime("vulkan-runtime-wrappers")])
+ tools.extend([add_runtime("mlir_vulkan_runtime")])
if config.enable_rocm_runner:
tools.extend([add_runtime("mlir_rocm_runtime")])
diff --git a/mlir/tools/CMakeLists.txt b/mlir/tools/CMakeLists.txt
index 072e83c5d45ea1..72a857b114fbf3 100644
--- a/mlir/tools/CMakeLists.txt
+++ b/mlir/tools/CMakeLists.txt
@@ -7,7 +7,6 @@ add_subdirectory(mlir-reduce)
add_subdirectory(mlir-rewrite)
add_subdirectory(mlir-shlib)
add_subdirectory(mlir-translate)
-add_subdirectory(mlir-vulkan-runner)
add_subdirectory(tblgen-lsp-server)
add_subdirectory(tblgen-to-irdl)
diff --git a/mlir/tools/mlir-vulkan-runner/CMakeLists.txt b/mlir/tools/mlir-vulkan-runner/CMakeLists.txt
deleted file mode 100644
index 26d6caacb0a7b1..00000000000000
--- a/mlir/tools/mlir-vulkan-runner/CMakeLists.txt
+++ /dev/null
@@ -1,99 +0,0 @@
-set(LLVM_OPTIONAL_SOURCES
- mlir-vulkan-runner.cpp
- vulkan-runtime-wrappers.cpp
- VulkanRuntime.cpp
- VulkanRuntime.h
- )
-
-if (MLIR_ENABLE_VULKAN_RUNNER)
- message(STATUS "Building the Vulkan runner")
-
- find_package(Vulkan)
-
- # If Vulkan is not found try a path specified by VULKAN_SDK.
- if (NOT Vulkan_FOUND)
- if ("$ENV{VULKAN_SDK}" STREQUAL "")
- message(FATAL_ERROR "Vulkan not found through CMake; please provide "
- "VULKAN_SDK path as an environment variable")
- endif()
-
- find_library(Vulkan_LIBRARY vulkan HINTS "$ENV{VULKAN_SDK}/lib" REQUIRED)
- if (Vulkan_LIBRARY)
- set(Vulkan_FOUND ON)
- set(Vulkan_INCLUDE_DIR "$ENV{VULKAN_SDK}/include")
- message(STATUS "Found Vulkan: " ${Vulkan_LIBRARY})
- endif()
- endif()
-
- if (NOT Vulkan_FOUND)
- message(FATAL_ERROR "Cannot find Vulkan library")
- endif()
-
- add_llvm_library(vulkan-runtime-wrappers SHARED
- vulkan-runtime-wrappers.cpp
- VulkanRuntime.cpp
- )
-
- target_include_directories(vulkan-runtime-wrappers
- PUBLIC
- ${Vulkan_INCLUDE_DIR}
- )
-
- # *IMPORTANT*: This library cannot depend on LLVM libraries. Otherwise,
- # it may cause LLVM version conflict when used together with other shared
- # libraries depending on LLVM. Notably, Mesa, who implements Vulkan
- # drivers on Linux, depends on the system libLLVM.so.
- target_link_libraries(vulkan-runtime-wrappers
- PUBLIC
- ${Vulkan_LIBRARY}
- )
-
- get_property(conversion_libs GLOBAL PROPERTY MLIR_CONVERSION_LIBS)
- set(LIBS
- ${conversion_libs}
- MLIRAnalysis
- MLIRArithDialect
- MLIRBuiltinToLLVMIRTranslation
- MLIRExecutionEngine
- MLIRFuncDialect
- MLIRGPUDialect
- MLIRIR
- MLIRJitRunner
- MLIRLLVMDialect
- MLIRLLVMCommonConversion
- MLIRLLVMToLLVMIRTranslation
- MLIRMemRefDialect
- MLIRMemRefToLLVM
- MLIRParser
- MLIRSPIRVDialect
- MLIRSPIRVTransforms
- MLIRSupport
- MLIRTargetLLVMIRExport
- MLIRTransforms
- MLIRTranslateLib
- MLIRVectorDialect
- MLIRVectorToLLVMPass
- ${Vulkan_LIBRARY}
- )
-
- # Manually expand the target library, since our MLIR libraries
- # aren't plugged into the LLVM dependency tracking. If we don't
- # do this then we can't insert the CodeGen library after ourselves
- llvm_expand_pseudo_components(TARGET_LIBS AllTargetsCodeGens)
- # Prepend LLVM in front of every target, this is how the library
- # are named with CMake
- SET(targets_to_link)
- FOREACH(t ${TARGET_LIBS})
- LIST(APPEND targets_to_link "LLVM${t}")
- ENDFOREACH(t)
-
- add_mlir_tool(mlir-vulkan-runner
- mlir-vulkan-runner.cpp
-
- DEPENDS
- vulkan-runtime-wrappers
- )
- llvm_update_compile_flags(mlir-vulkan-runner)
- target_link_libraries(mlir-vulkan-runner PRIVATE ${LIBS})
-
-endif()
diff --git a/mlir/tools/mlir-vulkan-runner/mlir-vulkan-runner.cpp b/mlir/tools/mlir-vulkan-runner/mlir-vulkan-runner.cpp
deleted file mode 100644
index 090df2d9ed2a5a..00000000000000
--- a/mlir/tools/mlir-vulkan-runner/mlir-vulkan-runner.cpp
+++ /dev/null
@@ -1,88 +0,0 @@
-//===- mlir-vulkan-runner.cpp - MLIR Vulkan Execution Driver --------------===//
-//
-// 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 is a command line utility that executes an MLIR file on the Vulkan by
-// translating MLIR GPU module to SPIR-V and host part to LLVM IR before
-// JIT-compiling and executing the latter.
-//
-//===----------------------------------------------------------------------===//
-
-#include "mlir/Conversion/ArithToLLVM/ArithToLLVM.h"
-#include "mlir/Conversion/ControlFlowToLLVM/ControlFlowToLLVM.h"
-#include "mlir/Conversion/FuncToLLVM/ConvertFuncToLLVMPass.h"
-#include "mlir/Conversion/GPUToVulkan/ConvertGPUToVulkanPass.h"
-#include "mlir/Conversion/LLVMCommon/LoweringOptions.h"
-#include "mlir/Conversion/MemRefToLLVM/MemRefToLLVM.h"
-#include "mlir/Conversion/ReconcileUnrealizedCasts/ReconcileUnrealizedCasts.h"
-#include "mlir/Conversion/VectorToLLVM/ConvertVectorToLLVMPass.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/LLVMIR/Transforms/RequestCWrappers.h"
-#include "mlir/Dialect/MemRef/IR/MemRef.h"
-#include "mlir/Dialect/MemRef/Transforms/Passes.h"
-#include "mlir/Dialect/SCF/IR/SCF.h"
-#include "mlir/Dialect/SPIRV/IR/SPIRVDialect.h"
-#include "mlir/Dialect/Vector/IR/VectorOps.h"
-#include "mlir/ExecutionEngine/JitRunner.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 "llvm/Support/InitLLVM.h"
-#include "llvm/Support/TargetSelect.h"
-
-using namespace mlir;
-
-static LogicalResult runMLIRPasses(Operation *op, JitRunnerOptions &) {
- auto module = dyn_cast<ModuleOp>(op);
- if (!module)
- return op->emitOpError("expected a 'builtin.module' op");
- PassManager passManager(module.getContext());
- if (failed(applyPassManagerCLOptions(passManager)))
- return failure();
-
- passManager.addPass(createConvertGpuLaunchFuncToVulkanLaunchFuncPass());
- passManager.addPass(createFinalizeMemRefToLLVMConversionPass());
- passManager.addPass(createConvertVectorToLLVMPass());
- passManager.nest<func::FuncOp>().addPass(LLVM::createRequestCWrappersPass());
- ConvertFuncToLLVMPassOptions funcToLLVMOptions{};
- funcToLLVMOptions.indexBitwidth =
- DataLayout(module).getTypeSizeInBits(IndexType::get(module.getContext()));
- passManager.addPass(createConvertFuncToLLVMPass(funcToLLVMOptions));
- passManager.addPass(createArithToLLVMConversionPass());
- passManager.addPass(createConvertControlFlowToLLVMPass());
- passManager.addPass(createReconcileUnrealizedCastsPass());
- passManager.addPass(createConvertVulkanLaunchFuncToVulkanCallsPass());
-
- return passManager.run(module);
-}
-
-int main(int argc, char **argv) {
- llvm::llvm_shutdown_obj x;
- registerPassManagerCLOptions();
-
- llvm::InitLLVM y(argc, argv);
- llvm::InitializeNativeTarget();
- llvm::InitializeNativeTargetAsmPrinter();
-
- mlir::JitRunnerConfig jitRunnerConfig;
- jitRunnerConfig.mlirTransformer = runMLIRPasses;
-
- mlir::DialectRegistry registry;
- registry.insert<mlir::arith::ArithDialect, mlir::LLVM::LLVMDialect,
- mlir::gpu::GPUDialect, mlir::spirv::SPIRVDialect,
- mlir::scf::SCFDialect, mlir::func::FuncDialect,
- mlir::memref::MemRefDialect, mlir::vector::VectorDialect>();
- mlir::registerBuiltinDialectTranslation(registry);
- mlir::registerLLVMDialectTranslation(registry);
-
- return mlir::JitRunnerMain(argc, argv, registry, jitRunnerConfig);
-}
More information about the Mlir-commits
mailing list