Google Git
Sign in
llvm/llvm-project/ba767d0bbbde4107700ff66ecfd97eb75d85a35d/./mlir/lib/Target/LLVM/XeVM/Target.cpp
blob: 28c654f3d7c9a54bf52850e8b4bbb4bb3670750d [file]
//===- Target.cpp - MLIR LLVM XeVM target compilation -----------*- 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
//
//===----------------------------------------------------------------------===//
//
// This files defines XeVM target related functions including registration
// calls for the `#xevm.target` compilation attribute.
//
//===----------------------------------------------------------------------===//
#include "mlir/Target/LLVM/XeVM/Target.h"
#include "mlir/Dialect/GPU/IR/CompilationInterfaces.h"
#include "mlir/Dialect/GPU/IR/GPUDialect.h"
#include "mlir/Dialect/LLVMIR/XeVMDialect.h"
#include "mlir/IR/BuiltinAttributeInterfaces.h"
#include "mlir/IR/BuiltinDialect.h"
#include "mlir/IR/BuiltinTypes.h"
#include "mlir/IR/DialectResourceBlobManager.h"
#include "mlir/Target/LLVM/XeVM/Utils.h"
#include "mlir/Target/LLVMIR/Dialect/GPU/GPUToLLVMIRTranslation.h"
#include "mlir/Target/LLVMIR/Dialect/LLVMIR/LLVMToLLVMIRTranslation.h"
#include "mlir/Target/LLVMIR/Dialect/XeVM/XeVMToLLVMIRTranslation.h"
#include "mlir/Target/LLVMIR/Export.h"
#include "llvm/ADT/ScopeExit.h"
#include "llvm/Bitcode/BitcodeWriter.h"
#include "llvm/Config/Targets.h"
#include "llvm/IR/LegacyPassManager.h"
#include "llvm/Support/FileSystem.h"
#include "llvm/Support/FileUtilities.h"
#include "llvm/Support/FormatVariadic.h"
#include "llvm/Support/MemoryBuffer.h"
#include "llvm/Support/Path.h"
#include "llvm/Support/Process.h"
#include "llvm/Support/Program.h"
#include "llvm/Support/TargetSelect.h"
#include "llvm/Support/raw_ostream.h"
#include "llvm/Target/TargetMachine.h"
#include <cstdint>
#include <cstdlib>
using namespace mlir;
using namespace mlir::xevm;
#if MLIR_XEVM_OCLOC_LIB_AVAILABLE
// Intel compute runtime includes libocloc in the distribution, but
// <ocloc_api.h> isn't included. Hence forward declarations for the libocloc
// shared-library APIs is needed. These replace the inclusion of <ocloc_api.h>
// so that the header is not a build-time requirement; the symbols are resolved
// at link/load time via the ocloc shared library.
extern "C" {
// Return code indicating successful ocloc compilation.
// Matches the OCLOC_SUCCESS enumerator in the real header (value 0).
enum OclocErrorCode : int { OCLOC_SUCCESS = 0 };
// Drives an in-process ocloc compilation.
/// Invokes ocloc API using C interface. Supported commands match
/// the functionality of ocloc executable (check ocloc's "help"
/// for reference : shared/offline_compiler/source/ocloc_api.cpp).
///
/// numArgs and argv params represent the command line.
/// Remaining params represent I/O.
/// Output params should be freed using oclocFreeOutput when
/// no longer needed.
/// List and names of outputs match outputs of ocloc executable.
///
/// \param numArgs is the number of arguments to pass to ocloc
///
/// \param argv is an array of arguments to be passed to ocloc
///
/// \param numSources is the number of in-memory representations
/// of source files to be passed to ocloc
///
/// \param dataSources is an array of in-memory representations
/// of source files to be passed to ocloc
///
/// \param lenSources is an array of sizes of in-memory representations
/// of source files passed to ocloc as dataSources
///
/// \param nameSources is an array of names of in-memory representations
/// of source files passed to ocloc as dataSources
///
/// \param numInputHeaders is the number of in-memory representations
/// of header files to be passed to ocloc
///
/// \param dataInputHeaders is an array of in-memory representations
/// of header files to be passed to ocloc
///
/// \param lenInputHeaders is an array of sizes of in-memory representations
/// of header files passed to ocloc as dataInputHeaders
///
/// \param nameInputHeaders is an array of names of in-memory representations
/// of header files passed to ocloc as dataInputHeaders
///
/// \param numOutputs returns the number of outputs
///
/// \param dataOutputs returns an array of in-memory representations
/// of output files
///
/// \param lenOutputs returns an array of sizes of in-memory representations
/// of output files
///
/// \param nameOutputs returns an array of names of in-memory representations
/// of output files. Special name stdout.log describes output that contains
/// messages generated by ocloc (e.g. compiler errors/warnings)
///
/// \returns 0 on success. Returns non-0 in case of failure.
int oclocInvoke(unsigned numArgs, const char **argv, unsigned numSources,
const uint8_t **dataSources, const uint64_t *lenSources,
const char **nameSources, unsigned numHeaders,
const uint8_t **dataHeaders, const uint64_t *lenHeaders,
const char **nameHeaders, unsigned *numOutputs,
uint8_t ***dataOutputs, uint64_t **lenOutputs,
char ***nameOutputs);
/// Frees results of oclocInvoke
///
/// \param numOutputs is number of outputs as returned by oclocInvoke
///
/// \param dataOutputs is array of outputs as returned by oclocInvoke
///
/// \param lenOutputs is array of sizes of outputs as returned by oclocInvoke
///
/// \param nameOutputs is array of names of outputs as returned by oclocInvoke
///
/// \returns 0 on success. Returns non-0 in case of failure.
int oclocFreeOutput(unsigned *numOutputs, uint8_t ***dataOutputs,
uint64_t **lenOutputs, char ***nameOutputs);
} // extern "C"
#endif // MLIR_XEVM_OCLOC_LIB_AVAILABLE
namespace {
// XeVM implementation of the gpu:TargetAttrInterface.
class XeVMTargetAttrImpl
: public gpu::TargetAttrInterface::FallbackModel<XeVMTargetAttrImpl> {
public:
std::optional<mlir::gpu::SerializedObject>
serializeToObject(Attribute attribute, Operation *module,
const gpu::TargetOptions &options) const;
Attribute createObject(Attribute attribute, Operation *module,
const mlir::gpu::SerializedObject &object,
const gpu::TargetOptions &options) const;
};
} // namespace
void mlir::xevm::registerXeVMTargetInterfaceExternalModels(
DialectRegistry &registry) {
registry.addExtension(+[](MLIRContext *ctx, XeVMDialect *dialect) {
XeVMTargetAttr::attachInterface<XeVMTargetAttrImpl>(*ctx);
});
}
void mlir::xevm::registerXeVMTargetInterfaceExternalModels(
MLIRContext &context) {
DialectRegistry registry;
registerXeVMTargetInterfaceExternalModels(registry);
context.appendDialectRegistry(registry);
}
SerializeGPUModuleBase::SerializeGPUModuleBase(
Operation &module, XeVMTargetAttr xeTarget,
const gpu::TargetOptions &targetOptions)
: ModuleToObject(module, xeTarget.getTriple(), "", {}, xeTarget.getO()),
xeTarget(xeTarget), librariesToLink(targetOptions.getLibrariesToLink()),
targetOptions(targetOptions) {
if (xeTarget.getLinkFiles())
librariesToLink.append(xeTarget.getLinkFiles().begin(),
xeTarget.getLinkFiles().end());
}
XeVMTargetAttr SerializeGPUModuleBase::getTarget() const { return xeTarget; }
std::optional<SmallVector<std::unique_ptr<llvm::Module>>>
SerializeGPUModuleBase::loadBitcodeFiles(llvm::Module &module) {
if (librariesToLink.empty())
return SmallVector<std::unique_ptr<llvm::Module>>();
SmallVector<std::unique_ptr<llvm::Module>> bcFiles;
if (failed(loadBitcodeFilesFromList(module.getContext(), librariesToLink,
bcFiles)))
return std::nullopt;
return std::move(bcFiles);
}
gpu::GPUModuleOp SerializeGPUModuleBase::getGPUModuleOp() {
return dyn_cast<gpu::GPUModuleOp>(&SerializeGPUModuleBase::getOperation());
}
// ----------------------------------------------------------------------------
// compile via the ocloc shared-library API (in-process, no temp files). Only
// compiled when the library is available at build time.
// ----------------------------------------------------------------------------
#if MLIR_XEVM_OCLOC_LIB_AVAILABLE
FailureOr<SmallVector<char, 0>>
SerializeGPUModuleBase::compileToBinaryViaLibocloc(StringRef asmStr,
StringRef inputFormat) {
Location loc = getGPUModuleOp().getLoc();
std::string asmFname = llvm::formatv(
"mlir-{0}-{1}-{2}.asm", getGPUModuleOp().getNameAttr().getValue(),
getTarget().getTriple(), getTarget().getChip());
// Build command-line options.
std::pair<llvm::BumpPtrAllocator, SmallVector<const char *>> cmdOpts =
targetOptions.tokenizeCmdOptions();
// Example: --gpu-module-to-binary="opts='opt1 opt2'"
const std::string cmdOptsStr = "\"" + llvm::join(cmdOpts.second, " ") + "\"";
std::vector<std::string> oclocArgs = {"ocloc",
"compile",
"-file",
asmFname,
inputFormat.str(),
"-device",
getTarget().getChip().str(),
"-options",
cmdOptsStr};
// Dump tool invocation commands.
#define DEBUG_TYPE "serialize-to-binary"
LLVM_DEBUG({
llvm::dbgs() << "libocloc invocation for module: "
<< getGPUModuleOp().getNameAttr() << "\n";
llvm::interleave(oclocArgs, llvm::dbgs(), " ");
llvm::dbgs() << "\n";
});
#undef DEBUG_TYPE
// Build a plain argv array expected by oclocInvoke.
std::vector<const char *> argv;
for (const auto &str : oclocArgs)
argv.push_back(str.c_str());
// Wire up in-memory source file.
const uint8_t *dataSources[1] = {
reinterpret_cast<const uint8_t *>(asmStr.data())};
const uint64_t lenSources[1] = {asmStr.size()};
const char *nameSources[1] = {asmFname.c_str()};
uint32_t outputsNum = 0;
uint8_t **outputs = nullptr;
uint64_t *outputLengths = nullptr;
char **outputNames = nullptr;
// Ensure ocloc output buffers are always freed on exit.
auto freeOutputs = llvm::scope_exit([&]() {
oclocFreeOutput(&outputsNum, &outputs, &outputLengths, &outputNames);
});
int err = oclocInvoke(static_cast<uint32_t>(argv.size()), argv.data(),
/*numSources=*/1, dataSources, lenSources, nameSources,
/*numHeaders=*/0, nullptr, nullptr, nullptr,
&outputsNum, &outputs, &outputLengths, &outputNames);
if (err != OCLOC_SUCCESS) {
emitError(loc) << "libocloc: `oclocInvoke` failed, error code: " << err;
// Emit any compiler log that ocloc produced.
for (uint32_t i = 0; i < outputsNum; ++i) {
if (llvm::StringRef(outputNames[i]).ends_with(".log"))
emitError(loc) << "Compiler log:\n"
<< llvm::StringRef(reinterpret_cast<char *>(outputs[i]),
outputLengths[i]);
}
return failure();
}
// Find and return the .bin output.
for (uint32_t i = 0; i < outputsNum; ++i) {
if (llvm::StringRef(outputNames[i]).ends_with(".bin")) {
char *begin = reinterpret_cast<char *>(outputs[i]);
return SmallVector<char, 0>(begin, begin + outputLengths[i]);
}
}
return emitError(loc) << "`oclocInvoke` did not produce `.bin` output";
}
#endif // MLIR_XEVM_OCLOC_LIB_AVAILABLE
// ----------------------------------------------------------------------------
// Compile by spawning the `ocloc` command-line tool as a process,
// communicating through temporary files. Acts as a fallback when the shared
// library is not available.
// ----------------------------------------------------------------------------
FailureOr<SmallVector<char, 0>>
SerializeGPUModuleBase::compileToBinaryViaOclocTool(StringRef asmStr,
StringRef inputFormat) {
using TmpFile = std::pair<llvm::SmallString<128>, llvm::FileRemover>;
// Locate the `ocloc` executable on PATH.
std::optional<std::string> oclocPath = findTool("ocloc");
if (!oclocPath) {
emitError(getGPUModuleOp().getLoc()) << "Could not find `ocloc` on PATH";
return failure();
}
Location loc = getGPUModuleOp().getLoc();
std::string basename = llvm::formatv(
"mlir-{0}-{1}-{2}", getGPUModuleOp().getNameAttr().getValue(),
getTarget().getTriple(), getTarget().getChip());
// Helper: create a named temporary file, returning path + auto-remover.
auto createTemp = [&](StringRef name,
StringRef suffix) -> std::optional<TmpFile> {
llvm::SmallString<128> path;
if (auto ec = llvm::sys::fs::createTemporaryFile(name, suffix, path)) {
emitError(loc) << "Couldn't create temp file `" << path
<< "`: " << ec.message();
return std::nullopt;
}
return TmpFile(path, llvm::FileRemover(path.c_str()));
};
std::optional<TmpFile> asmFile = createTemp(basename, "asm");
std::optional<TmpFile> binFile = createTemp(basename, "");
std::optional<TmpFile> logFile = createTemp(basename, "log");
if (!asmFile || !binFile || !logFile)
return failure();
// Write the assembly source to a temp file.
{
std::error_code ec;
llvm::raw_fd_ostream asmStream(asmFile->first, ec);
if (ec) {
emitError(loc) << "Couldn't open `" << asmFile->first
<< "`: " << ec.message();
return failure();
}
asmStream << asmStr;
if (asmStream.has_error()) {
emitError(loc) << "Error writing assembly to `" << asmFile->first << "`";
return failure();
}
asmStream.flush();
}
// Build command-line options.
std::pair<llvm::BumpPtrAllocator, SmallVector<const char *>> cmdOpts =
targetOptions.tokenizeCmdOptions();
const std::string cmdOptsStr = "\"" + llvm::join(cmdOpts.second, " ") + "\"";
SmallVector<StringRef, 12> oclocArgs(
{"ocloc", "compile", "-file", asmFile->first, inputFormat, "-device",
getTarget().getChip(), "-output", binFile->first, "-output_no_suffix",
"-options", cmdOptsStr});
// Dump tool invocation commands.
#define DEBUG_TYPE "serialize-to-binary"
LLVM_DEBUG({
llvm::dbgs() << "Tool invocation for module: "
<< getGPUModuleOp().getNameAttr() << "\n";
llvm::interleave(oclocArgs, llvm::dbgs(), " ");
llvm::dbgs() << "\n";
});
#undef DEBUG_TYPE
// Redirect stdout/stderr to the log temp file.
std::optional<StringRef> redirects[] = {std::nullopt, logFile->first,
logFile->first};
std::string errorMsg;
if (llvm::sys::ExecuteAndWait(*oclocPath, oclocArgs, std::nullopt, redirects,
0, 0, &errorMsg)) {
// Prefer a structured error message; otherwise dump the log file.
if (!errorMsg.empty()) {
emitError(loc) << "`ocloc` invocation failed: " << errorMsg;
} else if (auto log = llvm::MemoryBuffer::getFile(logFile->first)) {
emitError(loc) << "`ocloc` invocation failed. Log:\n"
<< (*log)->getBuffer();
} else {
emitError(loc) << "`ocloc` invocation failed (no log available)";
}
return failure();
}
// Read back the binary output (ocloc appends ".bin" to the base name).
binFile->first.append(".bin");
auto binaryBuffer = llvm::MemoryBuffer::getFile(binFile->first);
if (!binaryBuffer) {
emitError(loc) << "Couldn't open binary output `" << binFile->first
<< "`: " << binaryBuffer.getError().message();
return failure();
}
StringRef bin = (*binaryBuffer)->getBuffer();
return SmallVector<char, 0>(bin.begin(), bin.end());
}
// ----------------------------------------------------------------------------
// Public entry-point: prefer the in-process library path; fall back to the
// external tool when the library is not available.
// ----------------------------------------------------------------------------
FailureOr<SmallVector<char, 0>>
SerializeGPUModuleBase::compileToBinary(StringRef asmStr,
StringRef inputFormat) {
#if MLIR_XEVM_OCLOC_LIB_AVAILABLE
return compileToBinaryViaLibocloc(asmStr, inputFormat);
#else
return compileToBinaryViaOclocTool(asmStr, inputFormat);
#endif
}
std::optional<std::string> SerializeGPUModuleBase::findTool(StringRef tool) {
// 1. Check the toolkit path given in the command line.
StringRef pathRef = targetOptions.getToolkitPath();
SmallVector<char, 256> path;
if (!pathRef.empty()) {
path.insert(path.begin(), pathRef.begin(), pathRef.end());
llvm::sys::path::append(path, "bin", tool);
if (llvm::sys::fs::can_execute(path))
return StringRef(path.data(), path.size()).str();
}
// 2. Check PATH.
if (std::optional<std::string> toolPath =
llvm::sys::Process::FindInEnvPath("PATH", tool))
return *toolPath;
getGPUModuleOp().emitError()
<< "Couldn't find the `" << tool
<< "` binary. Please specify the toolkit "
"path via GpuModuleToBinaryPass or add the compiler to $PATH`.";
return std::nullopt;
}
namespace {
class SPIRVSerializer : public SerializeGPUModuleBase {
public:
SPIRVSerializer(Operation &module, XeVMTargetAttr xeTarget,
const gpu::TargetOptions &targetOptions)
: SerializeGPUModuleBase(module, xeTarget, targetOptions) {}
static void init();
/// Serializes the LLVM module to an object format, depending on the
/// compilation target selected in target options.
FailureOr<SmallVector<char, 0>>
moduleToObject(llvm::Module &llvmModule) override;
/// Runs the serialization pipeline, returning `std::nullopt` on error.
std::optional<SmallVector<char, 0>> run() override;
private:
/// Translates the LLVM module to SPIR-V binary using LLVM's
/// SPIR-V target.
std::optional<std::string>
translateToSPIRVBinary(llvm::Module &llvmModule,
llvm::TargetMachine &targetMachine);
};
} // namespace
void SPIRVSerializer::init() {
static llvm::once_flag initializeBackendOnce;
llvm::call_once(initializeBackendOnce, []() {
#if LLVM_HAS_SPIRV_TARGET
LLVMInitializeSPIRVTarget();
LLVMInitializeSPIRVTargetInfo();
LLVMInitializeSPIRVTargetMC();
LLVMInitializeSPIRVAsmPrinter();
#endif
});
}
#if LLVM_HAS_SPIRV_TARGET
static const std::vector<std::string> getDefaultSPIRVExtensions() {
return {
"SPV_EXT_relaxed_printf_string_address_space",
"SPV_INTEL_cache_controls",
"SPV_INTEL_variable_length_array",
};
}
namespace llvm {
class Module;
extern "C" bool
SPIRVTranslateModule(Module *M, std::string &SpirvObj, std::string &ErrMsg,
const std::vector<std::string> &AllowExtNames,
const std::vector<std::string> &Opts);
} // namespace llvm
#endif
// There is 1 way to finalize IL to native code: IGC
// There are 2 ways to access IGC: AOT (ocloc) and JIT (L0 runtime).
// - L0 runtime consumes IL and is external to MLIR codebase (rt wrappers).
// - `ocloc` tool can be "queried" from within MLIR.
FailureOr<SmallVector<char, 0>>
SPIRVSerializer::moduleToObject(llvm::Module &llvmModule) {
#define DEBUG_TYPE "serialize-to-llvm"
LLVM_DEBUG({
llvm::dbgs() << "LLVM IR for module: " << getGPUModuleOp().getNameAttr()
<< "\n";
llvm::dbgs() << llvmModule << "\n";
llvm::dbgs().flush();
});
#undef DEBUG_TYPE
// Return LLVM IR if the compilation target is `offload`.
if (targetOptions.getCompilationTarget() == gpu::CompilationTarget::Offload)
return SerializeGPUModuleBase::moduleToObject(llvmModule);
#if !LLVM_HAS_SPIRV_TARGET
return getGPUModuleOp()->emitError(
"The `SPIRV` target was not built. Please enable "
"it when building LLVM.");
#else
std::string serializedSPIRVBinary;
std::string ErrMsg;
std::vector<std::string> Opts;
Opts.push_back(triple.str());
Opts.push_back(std::to_string(optLevel));
// Translate the LLVM module to SPIR-V binary using LLVM's SPIR-V Backend API.
bool success =
SPIRVTranslateModule(&llvmModule, serializedSPIRVBinary, ErrMsg,
getDefaultSPIRVExtensions(), Opts);
if (!success)
return getGPUModuleOp().emitError()
<< "Failed translating the module to Binary."
<< "Error message: " << ErrMsg;
if (serializedSPIRVBinary.size() % 4)
return getGPUModuleOp().emitError()
<< "SPIRV code size must be a multiple of 4.";
StringRef spirvBin(serializedSPIRVBinary.c_str(),
serializedSPIRVBinary.size());
// Return SPIRV binary if the compilation target is `assembly`. Optimization
// and SPIR-V extensions are enabled for SPIR-V binary output in both paths
// (assembly and binary) as of now. SPIR-V binary
// is generated directly using the SPIR-V backends `SPIRVTranslateModule` API.
if (targetOptions.getCompilationTarget() ==
gpu::CompilationTarget::Assembly) {
#define DEBUG_TYPE "serialize-to-isa"
LLVM_DEBUG({
llvm::dbgs() << "SPIR-V for module: " << getGPUModuleOp().getNameAttr()
<< "\n";
llvm::dbgs() << serializedSPIRVBinary << "\n";
llvm::dbgs().flush();
});
#undef DEBUG_TYPE
return SmallVector<char, 0>(spirvBin.begin(), spirvBin.end());
}
// Return native binary. Compile the SPIR-V binary to native binary for Intel
// GPUs using `ocloc` compiler (Intel's OpenCL Offline Compiler).
return compileToBinary(spirvBin, "-spirv_input");
#endif // LLVM_HAS_SPIRV_TARGET
}
std::optional<SmallVector<char, 0>> SPIRVSerializer::run() {
// Translate the module to LLVM IR.
llvm::LLVMContext llvmContext;
std::unique_ptr<llvm::Module> llvmModule = translateToLLVMIR(llvmContext);
if (!llvmModule) {
getOperation().emitError() << "Failed creating the llvm::Module.";
return std::nullopt;
}
setDataLayoutAndTriple(*llvmModule);
if (initialLlvmIRCallback)
initialLlvmIRCallback(*llvmModule);
// Link bitcode files.
handleModulePreLink(*llvmModule);
{
auto libs = loadBitcodeFiles(*llvmModule);
if (!libs)
return std::nullopt;
if (!libs->empty())
if (failed(linkFiles(*llvmModule, std::move(*libs))))
return std::nullopt;
handleModulePostLink(*llvmModule);
}
if (linkedLlvmIRCallback)
linkedLlvmIRCallback(*llvmModule);
// Return the serialized object.
return moduleToObject(*llvmModule);
}
std::optional<std::string>
SPIRVSerializer::translateToSPIRVBinary(llvm::Module &llvmModule,
llvm::TargetMachine &targetMachine) {
std::string targetISA;
llvm::raw_string_ostream stream(targetISA);
{ // Drop pstream after this to prevent the ISA from being stuck buffering
llvm::buffer_ostream pstream(stream);
llvm::legacy::PassManager codegenPasses;
if (targetMachine.addPassesToEmitFile(codegenPasses, pstream, nullptr,
llvm::CodeGenFileType::ObjectFile))
return std::nullopt;
codegenPasses.run(llvmModule);
}
return targetISA;
}
std::optional<mlir::gpu::SerializedObject>
XeVMTargetAttrImpl::serializeToObject(Attribute attribute, Operation *module,
const gpu::TargetOptions &options) const {
if (!module)
return std::nullopt;
auto gpuMod = dyn_cast<gpu::GPUModuleOp>(module);
if (!gpuMod) {
module->emitError("expected to be a gpu.module op");
return std::nullopt;
}
auto xeTarget = cast<XeVMTargetAttr>(attribute);
if (xeTarget.getTriple().starts_with("spirv")) {
gpuMod.walk([&](LLVM::LLVMFuncOp funcOp) {
if (funcOp->hasAttr(gpu::GPUDialect::getKernelFuncAttrName())) {
funcOp.setIntelReqdSubGroupSize(16);
return WalkResult::interrupt();
}
return WalkResult::advance();
});
SPIRVSerializer serializer(*module, cast<XeVMTargetAttr>(attribute),
options);
serializer.init();
#if !LLVM_HAS_SPIRV_TARGET
module->emitError("Cannot run `TargetRegistry::lookupTarget()` for SPIRV "
"without having the target built.");
#endif
std::optional<SmallVector<char, 0>> binary = serializer.run();
if (!binary)
return std::nullopt;
return gpu::SerializedObject{std::move(*binary)};
}
module->emitError("Unsupported XeVM target triple: ") << xeTarget.getTriple();
return std::nullopt;
}
Attribute
XeVMTargetAttrImpl::createObject(Attribute attribute, Operation *module,
const mlir::gpu::SerializedObject &object,
const gpu::TargetOptions &options) const {
Builder builder(attribute.getContext());
gpu::CompilationTarget format = options.getCompilationTarget();
auto xeTarget = cast<XeVMTargetAttr>(attribute);
SmallVector<NamedAttribute, 2> properties;
if (format == gpu::CompilationTarget::Assembly)
properties.push_back(
builder.getNamedAttr("O", builder.getI32IntegerAttr(xeTarget.getO())));
DictionaryAttr objectProps;
if (!properties.empty())
objectProps = builder.getDictionaryAttr(properties);
return builder.getAttr<gpu::ObjectAttr>(
attribute, format,
builder.getStringAttr(
StringRef(object.getObject().data(), object.getObject().size())),
objectProps, /*kernels=*/nullptr);
}