================ @@ -0,0 +1,70 @@ +//===- SerializeToSPIRV.cpp - Convert GPU kernel to SPIRV blob -------------===// +// +// 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 +// +//===----------------------------------------------------------------------===// +/// +/// \file +/// This pass iterates all the SPIR-V modules in the top module and serializes +/// each SPIR-V module to SPIR-V binary and then attachs the binary blob as a +/// string attribute to the corresponding gpu module. +/// +//===----------------------------------------------------------------------===// + +#include "mlir/Dialect/GPU/Transforms/Passes.h" + +#include "mlir/Dialect/GPU/IR/GPUDialect.h" +#include "mlir/Dialect/GPU/Transforms/Passes.h" +#include "mlir/Dialect/SPIRV/IR/SPIRVDialect.h" +#include "mlir/Dialect/SPIRV/IR/SPIRVOps.h" +#include "mlir/Target/SPIRV/Serialization.h" + +namespace mlir { +#define GEN_PASS_DEF_GPUSERIALIZETOSPIRVPASS +#include "mlir/Dialect/GPU/Transforms/Passes.h.inc" +} // namespace mlir + +using namespace mlir; + +struct GpuSerializeToSPIRVPass : public mlir::impl::GpuSerializeToSPIRVPassBase<GpuSerializeToSPIRVPass> { +public: + void runOnOperation() override { + auto mod = getOperation(); + llvm::SmallVector<uint32_t, 0> spvBinary; + for (mlir::gpu::GPUModuleOp gpuMod : mod.getOps<gpu::GPUModuleOp>()) { + auto name = gpuMod.getName(); + // check that the spv module has the same name with gpu module except the + // prefix "__spv__" + auto isSameMod = [&](spirv::ModuleOp spvMod) -> bool { + auto spvModName = spvMod.getName(); + return spvModName->consume_front("__spv__") && spvModName == name; + }; + auto spvMods = mod.getOps<spirv::ModuleOp>(); + auto it = llvm::find_if(spvMods, isSameMod); + if (it == spvMods.end()) { + gpuMod.emitError() << "Unable to find corresponding SPIR-V module"; + signalPassFailure(); + return; + } + auto spvMod = *it; + + spvBinary.clear(); + // serialize the spv module to spv binary + if (mlir::failed(spirv::serialize(spvMod, spvBinary))) { + spvMod.emitError() << "Failed to serialize SPIR-V module"; + signalPassFailure(); + return; + } + + // attach the spv binary to the gpu module + auto spvData = + llvm::StringRef(reinterpret_cast<const char *>(spvBinary.data()), + spvBinary.size() * sizeof(uint32_t)); + auto spvAttr = mlir::StringAttr::get(&getContext(), spvData); + gpuMod->setAttr(gpu::getDefaultGpuBinaryAnnotation(), spvAttr); + spvMod->erase(); + } + } +}; ---------------- joker-eph wrote:
@fabianmcg : how would that fit in the new serialization flow? https://github.com/llvm/llvm-project/pull/65539 _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
