[Mlir-commits] [mlir] Add XeVM target and XeVM dialect integration tests. (PR #148286)
llvmlistbot at llvm.org
llvmlistbot at llvm.org
Fri Jul 11 13:17:33 PDT 2025
llvmbot wrote:
<!--LLVM PR SUMMARY COMMENT-->
@llvm/pr-subscribers-mlir-llvm
@llvm/pr-subscribers-mlir-gpu
Author: Sang Ik Lee (silee2)
<details>
<summary>Changes</summary>
Covers remaining parts required for XeVM dialect intgration testing. It has two high level components
- XeVM target and serialization support
- XeVM dialect integration tests using SYCL runtime
---
Patch is 59.90 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/148286.diff
21 Files Affected:
- (modified) mlir/CMakeLists.txt (+8)
- (modified) mlir/include/mlir/InitAllDialects.h (+2)
- (added) mlir/include/mlir/Target/LLVM/XeVM/Target.h (+31)
- (added) mlir/include/mlir/Target/LLVM/XeVM/Utils.h (+39)
- (modified) mlir/include/mlir/Target/LLVMIR/Dialect/All.h (+3)
- (added) mlir/include/mlir/Target/LLVMIR/Dialect/XeVM/XeVMToLLVMIRTranslation.h (+33)
- (modified) mlir/lib/Dialect/GPU/Transforms/XeVMAttachTarget.cpp (+1)
- (modified) mlir/lib/Target/LLVM/CMakeLists.txt (+32)
- (added) mlir/lib/Target/LLVM/XeVM/Target.cpp (+257)
- (modified) mlir/lib/Target/LLVMIR/CMakeLists.txt (+1)
- (modified) mlir/lib/Target/LLVMIR/Dialect/CMakeLists.txt (+1)
- (added) mlir/lib/Target/LLVMIR/Dialect/XeVM/CMakeLists.txt (+21)
- (added) mlir/lib/Target/LLVMIR/Dialect/XeVM/XeVMToLLVMIRTranslation.cpp (+117)
- (added) mlir/test/Integration/Dialect/XeVM/GPU/lit.local.cfg (+4)
- (added) mlir/test/Integration/Dialect/XeVM/GPU/xevm_block_dpas.mlir (+135)
- (added) mlir/test/Integration/Dialect/XeVM/GPU/xevm_block_load_store.mlir (+103)
- (added) mlir/test/Integration/Dialect/XeVM/GPU/xevm_block_load_store_pack_register.mlir (+119)
- (added) mlir/test/Integration/Dialect/XeVM/GPU/xevm_block_load_store_transpose.mlir (+127)
- (added) mlir/test/Integration/Dialect/XeVM/GPU/xevm_store_cst.mlir (+74)
- (modified) mlir/test/lib/Dialect/GPU/CMakeLists.txt (+1)
- (modified) mlir/test/lit.site.cfg.py.in (+1)
``````````diff
diff --git a/mlir/CMakeLists.txt b/mlir/CMakeLists.txt
index a1ad81f625cd6..7c9d62051f9f8 100644
--- a/mlir/CMakeLists.txt
+++ b/mlir/CMakeLists.txt
@@ -137,6 +137,14 @@ else()
set(MLIR_ENABLE_ROCM_CONVERSIONS 0)
endif()
+# Build the XeVM conversions and run according tests if the SPIRV backend
+# is available.
+if ("SPIRV" IN_LIST LLVM_TARGETS_TO_BUILD)
+ set(MLIR_ENABLE_XEVM_CONVERSIONS 1)
+else()
+ set(MLIR_ENABLE_XEVM_CONVERSIONS 0)
+endif()
+
set(MLIR_ENABLE_CUDA_RUNNER 0 CACHE BOOL "Enable building the MLIR CUDA runner")
set(MLIR_ENABLE_ROCM_RUNNER 0 CACHE BOOL "Enable building the MLIR ROCm runner")
set(MLIR_ENABLE_SYCL_RUNNER 0 CACHE BOOL "Enable building the MLIR SYCL runner")
diff --git a/mlir/include/mlir/InitAllDialects.h b/mlir/include/mlir/InitAllDialects.h
index c6fcf1a0d510b..79dcafe69f0a5 100644
--- a/mlir/include/mlir/InitAllDialects.h
+++ b/mlir/include/mlir/InitAllDialects.h
@@ -102,6 +102,7 @@
#include "mlir/Interfaces/CastInterfaces.h"
#include "mlir/Target/LLVM/NVVM/Target.h"
#include "mlir/Target/LLVM/ROCDL/Target.h"
+#include "mlir/Target/LLVM/XeVM/Target.h"
#include "mlir/Target/SPIRV/Target.h"
namespace mlir {
@@ -200,6 +201,7 @@ inline void registerAllDialects(DialectRegistry ®istry) {
NVVM::registerNVVMTargetInterfaceExternalModels(registry);
ROCDL::registerROCDLTargetInterfaceExternalModels(registry);
spirv::registerSPIRVTargetInterfaceExternalModels(registry);
+ xevm::registerXeVMTargetInterfaceExternalModels(registry);
}
/// Append all the MLIR dialects to the registry contained in the given context.
diff --git a/mlir/include/mlir/Target/LLVM/XeVM/Target.h b/mlir/include/mlir/Target/LLVM/XeVM/Target.h
new file mode 100644
index 0000000000000..31a93d0ebabfc
--- /dev/null
+++ b/mlir/include/mlir/Target/LLVM/XeVM/Target.h
@@ -0,0 +1,31 @@
+//===-- Target.h - MLIR XeVM target registration ----------------*- C++ -*-===//
+//
+// This file is licensed 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 provides registration calls for attaching the XeVM target interface.
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef MLIR_TARGET_XEVM_TARGET_H
+#define MLIR_TARGET_XEVM_TARGET_H
+
+namespace mlir {
+class DialectRegistry;
+class MLIRContext;
+} // namespace mlir
+
+namespace mlir::xevm {
+/// Registers the `TargetAttrInterface` for the `#xevm.target` attribute in
+/// the given registry.
+void registerXeVMTargetInterfaceExternalModels(mlir::DialectRegistry ®istry);
+
+/// Registers the `TargetAttrInterface` for the `#xevm.target` attribute in
+/// the registry associated with the given context.
+void registerXeVMTargetInterfaceExternalModels(mlir::MLIRContext &context);
+} // namespace mlir::xevm
+
+#endif // MLIR_TARGET_XEVM_TARGET_H
diff --git a/mlir/include/mlir/Target/LLVM/XeVM/Utils.h b/mlir/include/mlir/Target/LLVM/XeVM/Utils.h
new file mode 100644
index 0000000000000..c11a97f0d960a
--- /dev/null
+++ b/mlir/include/mlir/Target/LLVM/XeVM/Utils.h
@@ -0,0 +1,39 @@
+//===-- Utils.h - MLIR XeVM target utils ------------------------*- C++ -*-===//
+//
+// This file is licensed 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 declares XeVM target related utility classes and functions.
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef MLIR_TARGET_LLVM_XEVM_UTILS_H
+#define MLIR_TARGET_LLVM_XEVM_UTILS_H
+
+#include "mlir/Dialect/GPU/IR/CompilationInterfaces.h"
+#include "mlir/Dialect/LLVMIR/XeVMDialect.h"
+#include "mlir/Target/LLVM/ModuleToObject.h"
+
+namespace mlir {
+namespace xevm {
+
+/// Base class for all XeVM serializations from GPU modules into binary strings.
+/// By default this class serializes into LLVM bitcode.
+class SerializeGPUModuleBase : public mlir::LLVM::ModuleToObject {
+public:
+ SerializeGPUModuleBase(mlir::Operation &module, XeVMTargetAttr target,
+ const mlir::gpu::TargetOptions &targetOptions = {});
+
+ static void init();
+ XeVMTargetAttr getTarget() const;
+
+protected:
+ XeVMTargetAttr target;
+};
+} // namespace xevm
+} // namespace mlir
+
+#endif // MLIR_TARGET_LLVM_XEVM_UTILS_H
diff --git a/mlir/include/mlir/Target/LLVMIR/Dialect/All.h b/mlir/include/mlir/Target/LLVMIR/Dialect/All.h
index 60615cf601655..e4670cb1a9622 100644
--- a/mlir/include/mlir/Target/LLVMIR/Dialect/All.h
+++ b/mlir/include/mlir/Target/LLVMIR/Dialect/All.h
@@ -28,6 +28,7 @@
#include "mlir/Target/LLVMIR/Dialect/ROCDL/ROCDLToLLVMIRTranslation.h"
#include "mlir/Target/LLVMIR/Dialect/SPIRV/SPIRVToLLVMIRTranslation.h"
#include "mlir/Target/LLVMIR/Dialect/VCIX/VCIXToLLVMIRTranslation.h"
+#include "mlir/Target/LLVMIR/Dialect/XeVM/XeVMToLLVMIRTranslation.h"
namespace mlir {
class DialectRegistry;
@@ -47,6 +48,7 @@ static inline void registerAllToLLVMIRTranslations(DialectRegistry ®istry) {
registerROCDLDialectTranslation(registry);
registerSPIRVDialectTranslation(registry);
registerVCIXDialectTranslation(registry);
+ registerXeVMDialectTranslation(registry);
// Extension required for translating GPU offloading Ops.
gpu::registerOffloadingLLVMTranslationInterfaceExternalModels(registry);
@@ -63,6 +65,7 @@ registerAllGPUToLLVMIRTranslations(DialectRegistry ®istry) {
registerNVVMDialectTranslation(registry);
registerROCDLDialectTranslation(registry);
registerSPIRVDialectTranslation(registry);
+ registerXeVMDialectTranslation(registry);
// Extension required for translating GPU offloading Ops.
gpu::registerOffloadingLLVMTranslationInterfaceExternalModels(registry);
diff --git a/mlir/include/mlir/Target/LLVMIR/Dialect/XeVM/XeVMToLLVMIRTranslation.h b/mlir/include/mlir/Target/LLVMIR/Dialect/XeVM/XeVMToLLVMIRTranslation.h
new file mode 100644
index 0000000000000..149a2119657d5
--- /dev/null
+++ b/mlir/include/mlir/Target/LLVMIR/Dialect/XeVM/XeVMToLLVMIRTranslation.h
@@ -0,0 +1,33 @@
+//===-- XeVMToLLVMIRTranslation.h - XeVM to LLVM IR -------------*- C++ -*-===//
+//
+// This file is licensed 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 provides registration calls for XeVM dialect to LLVM IR translation.
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef MLIR_TARGET_LLVMIR_DIALECT_XEVM_XEVMTOLLVMIRTRANSLATION_H
+#define MLIR_TARGET_LLVMIR_DIALECT_XEVM_XEVMTOLLVMIRTRANSLATION_H
+
+namespace mlir {
+
+class DialectRegistry;
+class MLIRContext;
+} // namespace mlir
+
+namespace mlir {
+/// Register the XeVM dialect and the translation from it to the LLVM IR in the
+/// given registry;
+void registerXeVMDialectTranslation(mlir::DialectRegistry ®istry);
+
+/// Register the XeVM dialect and the translation from it in the registry
+/// associated with the given context.
+void registerXeVMDialectTranslation(mlir::MLIRContext &context);
+
+} // namespace mlir
+
+#endif // MLIR_TARGET_LLVMIR_DIALECT_XEVM_XEVMTOLLVMIRTRANSLATION_H
diff --git a/mlir/lib/Dialect/GPU/Transforms/XeVMAttachTarget.cpp b/mlir/lib/Dialect/GPU/Transforms/XeVMAttachTarget.cpp
index e9cf4939a13b8..6da76e9e7a331 100644
--- a/mlir/lib/Dialect/GPU/Transforms/XeVMAttachTarget.cpp
+++ b/mlir/lib/Dialect/GPU/Transforms/XeVMAttachTarget.cpp
@@ -17,6 +17,7 @@
#include "mlir/Dialect/LLVMIR/XeVMDialect.h"
#include "mlir/IR/Builders.h"
#include "mlir/Pass/Pass.h"
+#include "mlir/Target/LLVM/XeVM/Target.h"
#include "llvm/Support/Regex.h"
namespace mlir {
diff --git a/mlir/lib/Target/LLVM/CMakeLists.txt b/mlir/lib/Target/LLVM/CMakeLists.txt
index 83fbf7a5fe5f3..ed15c5d2ab2ca 100644
--- a/mlir/lib/Target/LLVM/CMakeLists.txt
+++ b/mlir/lib/Target/LLVM/CMakeLists.txt
@@ -209,3 +209,35 @@ if(MLIR_ENABLE_ROCM_CONVERSIONS)
)
endif()
+if ("SPIRV" IN_LIST LLVM_TARGETS_TO_BUILD)
+ set(SPIRV_LIBS
+ SPIRVCodeGen
+
+ )
+endif()
+
+add_mlir_dialect_library(MLIRXeVMTarget
+ XeVM/Target.cpp
+
+ OBJECT
+
+ ADDITIONAL_HEADER_DIRS
+ ${MLIR_MAIN_INCLUDE_DIR}/mlir/Dialect/LLVMIR
+
+ LINK_COMPONENTS
+ ${SPIRV_LIBS}
+
+ LINK_LIBS PUBLIC
+ MLIRIR
+ MLIRExecutionEngineUtils
+ MLIRSupport
+ MLIRGPUDialect
+ MLIRTargetLLVM
+ MLIRXeVMToLLVMIRTranslation
+)
+
+# Ensure SPIRV headers are included. Warning: references build directory!
+target_include_directories(MLIRXeVMTarget PRIVATE
+ ${LLVM_MAIN_SRC_DIR}/lib/Target/SPIRV
+ ${LLVM_BINARY_DIR}/lib/Target/SPIRV
+)
diff --git a/mlir/lib/Target/LLVM/XeVM/Target.cpp b/mlir/lib/Target/LLVM/XeVM/Target.cpp
new file mode 100644
index 0000000000000..380e2bff222ca
--- /dev/null
+++ b/mlir/lib/Target/LLVM/XeVM/Target.cpp
@@ -0,0 +1,257 @@
+//===-- Target.cpp - MLIR LLVM XeVM target compilation ----------*- C++ -*-===//
+//
+// This file is licensed 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 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/LLVMDialect.h"
+#include "mlir/Dialect/LLVMIR/XeVMDialect.h"
+#include "mlir/IR/ExtensibleDialect.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 "llvm/Config/Targets.h"
+#include "llvm/IR/LegacyPassManager.h"
+#include "llvm/Support/MemoryBuffer.h"
+#include "llvm/Support/TargetSelect.h"
+#include "llvm/Target/TargetMachine.h"
+
+#include "llvm/IR/Function.h"
+#include "llvm/IR/IRBuilder.h"
+#include "llvm/IR/Instructions.h"
+#include "llvm/IR/Module.h"
+
+// FIXME: One of the headers uses `.inc` file from the build directory, this
+// does not work for installation (i.e., DCMAKE_INSTALL_PREFIX) caching as build
+// directory will not be cached. Since float atomics are not yet supported by
+// the backend anyway, we can afford to temporarily comment this section.
+
+// #if LLVM_HAS_SPIRV_TARGET
+// #pragma GCC diagnostic push
+// #pragma GCC diagnostic ignored "-Wnon-virtual-dtor"
+// #include "SPIRVTargetMachine.h"
+// #pragma GCC diagnostic pop
+
+// #include "SPIRVCommandLine.h"
+// #endif // LLVM_HAS_SPIRV_TARGET
+
+#include <set>
+
+using namespace mlir;
+
+namespace {
+// XeVM implementation of the gpu:TargetAttrInterface.
+class XeVMTargetAttrImpl
+ : public gpu::TargetAttrInterface::FallbackModel<XeVMTargetAttrImpl> {
+public:
+ std::optional<SmallVector<char, 0>>
+ serializeToObject(Attribute attribute, Operation *module,
+ const gpu::TargetOptions &options) const;
+
+ Attribute createObject(Attribute attribute, Operation *module,
+ const SmallVector<char, 0> &object,
+ const gpu::TargetOptions &options) const;
+};
+} // namespace
+
+void mlir::xevm::registerXeVMTargetInterfaceExternalModels(
+ DialectRegistry ®istry) {
+ registry.addExtension(
+ +[](MLIRContext *ctx, mlir::xevm::XeVMDialect *dialect) {
+ mlir::xevm::XeVMTargetAttr::attachInterface<XeVMTargetAttrImpl>(*ctx);
+ });
+}
+
+void mlir::xevm::registerXeVMTargetInterfaceExternalModels(
+ MLIRContext &context) {
+ DialectRegistry registry;
+ registerXeVMTargetInterfaceExternalModels(registry);
+ context.appendDialectRegistry(registry);
+}
+
+mlir::xevm::SerializeGPUModuleBase::SerializeGPUModuleBase(
+ Operation &module, mlir::xevm::XeVMTargetAttr target,
+ const gpu::TargetOptions &targetOptions)
+ : ModuleToObject(module, target.getTriple(), "", {}, target.getO()),
+ target(target) {}
+
+void mlir::xevm::SerializeGPUModuleBase::init() {
+ static llvm::once_flag initializeBackendOnce;
+ llvm::call_once(initializeBackendOnce, []() {
+#if LLVM_HAS_SPIRV_TARGET
+ LLVMInitializeSPIRVTarget();
+ LLVMInitializeSPIRVTargetInfo();
+ LLVMInitializeSPIRVTargetMC();
+ LLVMInitializeSPIRVAsmPrinter();
+#endif
+ });
+}
+
+mlir::xevm::XeVMTargetAttr
+mlir::xevm::SerializeGPUModuleBase::getTarget() const {
+ return target;
+}
+
+namespace {
+class SpirSerializer : public mlir::xevm::SerializeGPUModuleBase {
+public:
+ SpirSerializer(Operation &module, mlir::xevm::XeVMTargetAttr target,
+ const gpu::TargetOptions &targetOptions)
+ : mlir::xevm::SerializeGPUModuleBase(module, target, targetOptions) {}
+
+ gpu::GPUModuleOp getOperation();
+
+ std::optional<SmallVector<char, 0>>
+ moduleToObject(llvm::Module &llvmModule) override;
+
+private:
+ std::optional<std::string>
+ translateToSPIRVBinary(llvm::Module &llvmModule,
+ llvm::TargetMachine &targetMachine);
+ gpu::TargetOptions targetOptions;
+};
+} // namespace
+
+gpu::GPUModuleOp SpirSerializer::getOperation() {
+ return dyn_cast<gpu::GPUModuleOp>(
+ &mlir::xevm::SerializeGPUModuleBase::getOperation());
+}
+
+std::optional<SmallVector<char, 0>>
+SpirSerializer::moduleToObject(llvm::Module &llvmModule) {
+ // Return LLVM IR if the compilation target is `offload`.
+ if (targetOptions.getCompilationTarget() == gpu::CompilationTarget::Offload)
+ return mlir::xevm::SerializeGPUModuleBase::moduleToObject(llvmModule);
+
+#if !LLVM_HAS_SPIRV_TARGET
+ getOperation()->emitError(
+ "The `SPIRV` target was not built. Please enable it when building LLVM.");
+ return std::nullopt;
+#endif // LLVM_HAS_SPIRV_TARGET
+
+ std::optional<llvm::TargetMachine *> targetMachine =
+ getOrCreateTargetMachine();
+ if (!targetMachine) {
+ getOperation().emitError() << "Target Machine unavailable for triple "
+ << triple << ", can't compile with LLVM\n";
+ return std::nullopt;
+ }
+
+ //===----------------------------------------------------------------------===//
+ // Workaround to enable spirv extensions that are not added to target machine
+ // by default.
+
+ // FIXME: see fixme comment above SPIRV headers.
+ // #if LLVM_HAS_SPIRV_TARGET
+ // std::set<llvm::SPIRV::Extension::Extension> AllowedExtIds{
+ // llvm::SPIRV::Extension::Extension::SPV_EXT_shader_atomic_float_add,
+ // llvm::SPIRV::Extension::Extension::SPV_EXT_shader_atomic_float16_add};
+ // llvm::SPIRVTargetMachine *STM =
+ // static_cast<llvm::SPIRVTargetMachine *>(targetMachine.value());
+ // const_cast<llvm::SPIRVSubtarget *>(STM->getSubtargetImpl())
+ // ->initAvailableExtensions(AllowedExtIds);
+ // #endif // LLVM_HAS_SPIRV_TARGET
+
+ //===----------------------------------------------------------------------===//
+
+ // Return SPIRV if the compilation target is `assembly`.
+ if (targetOptions.getCompilationTarget() ==
+ gpu::CompilationTarget::Assembly) {
+ std::optional<std::string> serializedISA =
+ translateToISA(llvmModule, **targetMachine);
+ if (!serializedISA) {
+ getOperation().emitError() << "Failed translating the module to ISA.";
+ return std::nullopt;
+ }
+ // Make sure to include the null terminator.
+ StringRef bin(serializedISA->c_str(), serializedISA->size() + 1);
+ return SmallVector<char, 0>(bin.begin(), bin.end());
+ }
+
+ std::optional<std::string> serializedSPIRVBinary =
+ translateToSPIRVBinary(llvmModule, **targetMachine);
+ if (!serializedSPIRVBinary) {
+ getOperation().emitError() << "Failed translating the module to Binary.";
+ return std::nullopt;
+ }
+ if (serializedSPIRVBinary->size() % 4) {
+ getOperation().emitError() << "SPIRV code size must be a multiple of 4.";
+ return std::nullopt;
+ }
+ StringRef bin(serializedSPIRVBinary->c_str(), serializedSPIRVBinary->size());
+ return SmallVector<char, 0>(bin.begin(), bin.end());
+}
+
+std::optional<std::string>
+SpirSerializer::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<SmallVector<char, 0>>
+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;
+ }
+ gpuMod.walk([&](LLVM::LLVMFuncOp funcOp) {
+ if (funcOp->hasAttr(gpu::GPUDialect::getKernelFuncAttrName())) {
+ funcOp.setIntelReqdSubGroupSize(16);
+ return WalkResult::interrupt();
+ }
+ return WalkResult::advance();
+ });
+
+ SpirSerializer serializer(
+ *module, cast<mlir::xevm::XeVMTargetAttr>(attribute), options);
+ serializer.init();
+
+#if !LLVM_HAS_SPIRV_TARGET
+ module->emitError("Cannot run `TargetRegistry::lookupTarget()` for SPIRV "
+ "without having the target built.");
+#endif
+
+ return serializer.run();
+}
+
+Attribute
+XeVMTargetAttrImpl::createObject(Attribute attribute, Operation *module,
+ const SmallVector<char, 0> &object,
+ const gpu::TargetOptions &options) const {
+ gpu::CompilationTarget format = options.getCompilationTarget();
+ DictionaryAttr objectProps;
+ Builder builder(attribute.getContext());
+ return builder.getAttr<gpu::ObjectAttr>(
+ attribute, format,
+ builder.getStringAttr(StringRef(object.data(), object.size())),
+ objectProps, /*kernels=*/nullptr);
+}
diff --git a/mlir/lib/Target/LLVMIR/CMakeLists.txt b/mlir/lib/Target/LLVMIR/CMakeLists.txt
index af22a7ff04bf0..9ea5c6835e8ef 100644
--- a/mlir/lib/Target/LLVMIR/CMakeLists.txt
+++ b/mlir/lib/Target/LLVMIR/CMakeLists.txt
@@ -60,6 +60,7 @@ add_mlir_translation_library(MLIRToLLVMIRTranslationRegistration
MLIRROCDLToLLVMIRTranslation
MLIRSPIRVToLLVMIRTranslation
MLIRVCIXToLLVMIRTranslation
+ MLIRXeVMToLLVMIRTranslation
)
add_mlir_translation_library(MLIRTargetLLVMIRImport
diff --git a/mlir/lib/Target/LLVMIR/Dialect/CMakeLists.txt b/mlir/lib/Target/LLVMIR/Dialect/CMakeLists.txt
index f030fa78942d5..86c731a1074c3 100644
--- a/mlir/lib/Target/LLVMIR/Dialect/CMakeLists.txt
+++ b/mlir/lib/Target/LLVMIR/Dialect/CMakeLists.txt
@@ -10,3 +10,4 @@ add_subdirectory(OpenMP)
add_subdirectory(ROCDL)
add_subdirectory(SPIRV)
add_subdirectory(VCIX)
+add_subdirectory(XeVM)
diff --git a/mlir/lib/Target/LLVMIR/Dialect/XeVM/CMakeLists.txt b/mlir/lib/Target/LLVMIR/Dialect/XeVM/CMakeLists.txt
new file mode 100644
index 0000000000000..6308d7e2e4404
--- /dev/null
+++ b/mlir/lib/Target/LLVMIR/Dialect/XeVM/CMakeLists.txt
@@ -0,0 +1,21 @@
+set(LLVM_OPTIONAL_SOURCES
+ XeVMToLLVMIRTranslation.cpp
+)
+
+add_mlir_translation_library(MLIRXeVMToLLVMIRTranslation
+ XeVMToLLVMIRTranslation.cpp
+
+ DEPENDS
+ MLIRXeVMConversionsIncGen
+
+ LINK_COMPONENTS
+ Core
+
+ LINK_LIBS PUBLIC
+ MLIRDialectUtils
+ MLIRIR
+ MLIRLLVMDialect
+ MLIRXeVMDialect
+ MLIRSupport
+ MLIRTargetLLVMIRExport
+)
diff --git a/mlir/lib/Target/LLVMIR/Dialect/XeVM/XeVMToLLVMIRTranslation.cpp b/mlir/lib/Target/LLVMIR/Dialect/XeVM/XeVMToLLVMIRTran...
[truncated]
``````````
</details>
https://github.com/llvm/llvm-project/pull/148286
More information about the Mlir-commits
mailing list