TheRock/compiler/amd-llvm/mlir/lib/Target/LLVM/XeVM/Target.cpp

//===- 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
// Enable default set of SPIR-V extensions.
// Supported extensions vary based on chip. See,
// https://github.com/intel/intel-graphics-compiler/blob/master/
//   documentation/igc/supported-spirv-extensions.md
// for details.
static const std::vector<std::string> getDefaultSPIRVExtensions() {
  return {
      "SPV_KHR_bfloat16",
      "SPV_EXT_relaxed_printf_string_address_space",
      "SPV_INTEL_cache_controls",
      "SPV_INTEL_variable_length_array",
      "SPV_INTEL_bfloat16_arithmetic",
  };
}

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