caio/clang-p2996
1
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

[mlir] Remove mlir-vulkan-runner and GPUToVulkan conversion passes (#123750)

Browse Source 
This follows up on 733be4ed7d, 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.
This commit is contained in:
Andrea Faulds
2025-01-21 16:51:27 +01:00
committed by GitHub
parent 5deb4ef9ab
commit e7e3c45bc7
33 changed files with 95 additions and 1158 deletions
Download Patch File Download Diff File Expand all files Collapse all files

36
mlir/include/mlir/Conversion/GPUToVulkan/ConvertGPUToVulkanPass.h
View File

@@ -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

1
mlir/include/mlir/Conversion/Passes.h
View File

@@ -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"

25
mlir/include/mlir/Conversion/Passes.td
View File

@@ -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
//===----------------------------------------------------------------------===//

1
mlir/lib/Conversion/CMakeLists.txt
View File

@@ -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)

19
mlir/lib/Conversion/GPUToVulkan/CMakeLists.txt
View File

@@ -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
)

219
mlir/lib/Conversion/GPUToVulkan/ConvertGPULaunchFuncToVulkanLaunchFunc.cpp
View File

@@ -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>();
}

448
mlir/lib/Conversion/GPUToVulkan/ConvertLaunchFuncToVulkanCalls.cpp
View File

@@ -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();
}

45
mlir/lib/ExecutionEngine/CMakeLists.txt
View File

@@ -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()

0
mlir/tools/mlir-vulkan-runner/VulkanRuntime.cpp → mlir/lib/ExecutionEngine/VulkanRuntime.cpp
View File

0
mlir/tools/mlir-vulkan-runner/VulkanRuntime.h → mlir/lib/ExecutionEngine/VulkanRuntime.h
View File

77
mlir/tools/mlir-vulkan-runner/vulkan-runtime-wrappers.cpp → mlir/lib/ExecutionEngine/VulkanRuntimeWrappers.cpp
View File

@@ -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

2
mlir/test/CMakeLists.txt
View File

@@ -206,7 +206,7 @@ endif()
if(MLIR_ENABLE_VULKAN_RUNNER)
list(APPEND MLIR_TEST_DEPENDS
mlir-vulkan-runner
mlir_vulkan_runtime
)
endif()

62
mlir/test/Conversion/GPUToVulkan/invoke-vulkan.mlir
View File

@@ -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)
}

35
mlir/test/Conversion/GPUToVulkan/lower-gpu-launch-vulkan-launch.mlir
View File

@@ -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
}
}

4
mlir/test/mlir-vulkan-runner/addf.mlir → mlir/test/Integration/GPU/Vulkan/addf.mlir
View File

@@ -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 {

4
mlir/test/mlir-vulkan-runner/addf_if.mlir → mlir/test/Integration/GPU/Vulkan/addf_if.mlir
View File

@@ -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 {

4
mlir/test/mlir-vulkan-runner/addi.mlir → mlir/test/Integration/GPU/Vulkan/addi.mlir
View File

@@ -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 {

4
mlir/test/mlir-vulkan-runner/addi8.mlir → mlir/test/Integration/GPU/Vulkan/addi8.mlir
View File

@@ -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 {

8
mlir/test/mlir-vulkan-runner/addui_extended.mlir → mlir/test/Integration/GPU/Vulkan/addui_extended.mlir
View File

@@ -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]

0
mlir/test/mlir-vulkan-runner/lit.local.cfg → mlir/test/Integration/GPU/Vulkan/lit.local.cfg
View File

4
mlir/test/mlir-vulkan-runner/mulf.mlir → mlir/test/Integration/GPU/Vulkan/mulf.mlir
View File

@@ -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 {

8
mlir/test/mlir-vulkan-runner/smul_extended.mlir → mlir/test/Integration/GPU/Vulkan/smul_extended.mlir
View File

@@ -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]

4
mlir/test/mlir-vulkan-runner/subf.mlir → mlir/test/Integration/GPU/Vulkan/subf.mlir
View File

@@ -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 {

4
mlir/test/mlir-vulkan-runner/time.mlir → mlir/test/Integration/GPU/Vulkan/time.mlir
View File

@@ -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

8
mlir/test/mlir-vulkan-runner/umul_extended.mlir → mlir/test/Integration/GPU/Vulkan/umul_extended.mlir
View File

@@ -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]

4
mlir/test/mlir-vulkan-runner/vector-deinterleave.mlir → mlir/test/Integration/GPU/Vulkan/vector-deinterleave.mlir
View File

@@ -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]

4
mlir/test/mlir-vulkan-runner/vector-interleave.mlir → mlir/test/Integration/GPU/Vulkan/vector-interleave.mlir
View File

@@ -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]

4
mlir/test/mlir-vulkan-runner/vector-shuffle.mlir → mlir/test/Integration/GPU/Vulkan/vector-shuffle.mlir
View File

@@ -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]

29
mlir/test/lib/Pass/TestVulkanRunnerPipeline.cpp
View File

@@ -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

2
mlir/test/lit.cfg.py
View File

@@ -197,7 +197,7 @@ tools = [
]
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")])

1
mlir/tools/CMakeLists.txt
View File

@@ -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)

99
mlir/tools/mlir-vulkan-runner/CMakeLists.txt
View File

@@ -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()

88
mlir/tools/mlir-vulkan-runner/mlir-vulkan-runner.cpp
View File

@@ -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);
}