[Mlir-commits] [mlir] [MLIR] Add SyclRuntimeWrapper (PR #69648)
Nishant Patel
llvmlistbot at llvm.org
Thu Oct 26 08:56:54 PDT 2023
https://github.com/nbpatel updated https://github.com/llvm/llvm-project/pull/69648
>From 843d9b200f3708676119ed0b2b3b0657cd3fd47b Mon Sep 17 00:00:00 2001
From: Nishant Patel <nishant.b.patel at intel.com>
Date: Mon, 18 Sep 2023 18:26:22 +0000
Subject: [PATCH 1/6] Add SYCL runtimet wrappers
---
.../ExecutionEngine/SyclRuntimeWrappers.cpp | 221 ++++++++++++++++++
1 file changed, 221 insertions(+)
create mode 100644 mlir/lib/ExecutionEngine/SyclRuntimeWrappers.cpp
diff --git a/mlir/lib/ExecutionEngine/SyclRuntimeWrappers.cpp b/mlir/lib/ExecutionEngine/SyclRuntimeWrappers.cpp
new file mode 100644
index 000000000000000..5bb58ea0dbe0c21
--- /dev/null
+++ b/mlir/lib/ExecutionEngine/SyclRuntimeWrappers.cpp
@@ -0,0 +1,221 @@
+//===- SyclRuntimeWrappers.cpp - MLIR SYCL 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.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+//
+// Implements C wrappers around the sycl runtime library.
+//
+//===----------------------------------------------------------------------===//
+
+#include <algorithm>
+#include <array>
+#include <atomic>
+#include <cassert>
+#include <cfloat>
+#include <cstdint>
+#include <cstdio>
+#include <cstdlib>
+#include <stdexcept>
+#include <tuple>
+#include <vector>
+
+#include <CL/sycl.hpp>
+#include <level_zero/ze_api.h>
+#include <map>
+#include <mutex>
+#include <sycl/ext/oneapi/backend/level_zero.hpp>
+
+#ifdef _WIN32
+#define SYCL_RUNTIME_EXPORT __declspec(dllexport)
+#else
+#define SYCL_RUNTIME_EXPORT
+#endif // _WIN32
+
+namespace {
+
+template <typename F>
+auto catchAll(F &&func) {
+ try {
+ return func();
+ } catch (const std::exception &e) {
+ fprintf(stdout, "An exception was thrown: %s\n", e.what());
+ fflush(stdout);
+ abort();
+ } catch (...) {
+ fprintf(stdout, "An unknown exception was thrown\n");
+ fflush(stdout);
+ abort();
+ }
+}
+
+#define L0_SAFE_CALL(call) \
+ { \
+ ze_result_t status = (call); \
+ if (status != ZE_RESULT_SUCCESS) { \
+ fprintf(stdout, "L0 error %d\n", status); \
+ fflush(stdout); \
+ abort(); \
+ } \
+ }
+
+} // namespace
+
+static sycl::device getDefaultDevice() {
+ auto platformList = sycl::platform::get_platforms();
+ for (const auto &platform : platformList) {
+ auto platformName = platform.get_info<sycl::info::platform::name>();
+ bool isLevelZero = platformName.find("Level-Zero") != std::string::npos;
+ if (!isLevelZero)
+ continue;
+
+ return platform.get_devices()[0];
+ }
+ throw std::runtime_error("getDefaultDevice failed");
+}
+
+#pragma clang diagnostic push
+#pragma clang diagnostic ignored "-Wglobal-constructors"
+
+// Create global device and context
+sycl::device syclDevice = getDefaultDevice();
+sycl::context syclContext { syclDevice };
+
+#pragma clang diagnostic pop
+
+struct QUEUE {
+ sycl::queue syclQueue_;
+
+ QUEUE() { syclQueue_ = sycl::queue(syclContext, syclDevice); }
+};
+
+static void *allocDeviceMemory(QUEUE *queue, size_t size, bool isShared) {
+ void *memPtr = nullptr;
+ if (isShared) {
+ memPtr = sycl::aligned_alloc_shared(64, size, syclDevice, syclContext);
+ } else {
+ memPtr = sycl::aligned_alloc_device(64, size, syclDevice, syclContext);
+ }
+ if (memPtr == nullptr) {
+ throw std::runtime_error("mem allocation failed!");
+ }
+ return memPtr;
+}
+
+static void deallocDeviceMemory(QUEUE *queue, void *ptr) {
+ sycl::free(ptr, queue->syclQueue_);
+}
+
+static ze_module_handle_t loadModule(const void *data, size_t dataSize) {
+ assert(data);
+ ze_module_handle_t zeModule;
+ ze_module_desc_t desc = {ZE_STRUCTURE_TYPE_MODULE_DESC,
+ nullptr,
+ ZE_MODULE_FORMAT_IL_SPIRV,
+ dataSize,
+ (const uint8_t *)data,
+ nullptr,
+ nullptr};
+ auto zeDevice =
+ sycl::get_native<sycl::backend::ext_oneapi_level_zero>(syclDevice);
+ auto zeContext =
+ sycl::get_native<sycl::backend::ext_oneapi_level_zero>(syclContext);
+ L0_SAFE_CALL(zeModuleCreate(zeContext, zeDevice, &desc, &zeModule, nullptr));
+ return zeModule;
+}
+
+static sycl::kernel *getKernel(ze_module_handle_t zeModule, const char *name) {
+ assert(zeModule);
+ assert(name);
+ ze_kernel_handle_t zeKernel;
+ sycl::kernel *syclKernel;
+ ze_kernel_desc_t desc = {};
+ desc.pKernelName = name;
+
+ L0_SAFE_CALL(zeKernelCreate(zeModule, &desc, &zeKernel));
+ sycl::kernel_bundle<sycl::bundle_state::executable> kernelBundle =
+ sycl::make_kernel_bundle<sycl::backend::ext_oneapi_level_zero,
+ sycl::bundle_state::executable>({zeModule},
+ syclContext);
+
+ auto kernel = sycl::make_kernel<sycl::backend::ext_oneapi_level_zero>(
+ {kernelBundle, zeKernel}, syclContext);
+ syclKernel = new sycl::kernel(kernel);
+ return syclKernel;
+}
+
+static void launchKernel(QUEUE *queue, sycl::kernel *kernel, size_t gridX,
+ size_t gridY, size_t gridZ, size_t blockX,
+ size_t blockY, size_t blockZ, size_t sharedMemBytes,
+ void **params, size_t paramsCount) {
+ auto syclGlobalRange =
+ ::sycl::range<3>(blockZ * gridZ, blockY * gridY, blockX * gridX);
+ auto syclLocalRange = ::sycl::range<3>(blockZ, blockY, blockX);
+ sycl::nd_range<3> syclNdRange(
+ sycl::nd_range<3>(syclGlobalRange, syclLocalRange));
+
+ queue->syclQueue_.submit([&](sycl::handler &cgh) {
+ for (size_t i = 0; i < paramsCount; i++) {
+ cgh.set_arg(static_cast<uint32_t>(i), *(static_cast<void **>(params[i])));
+ }
+ cgh.parallel_for(syclNdRange, *kernel);
+ });
+}
+
+extern "C" SYCL_RUNTIME_EXPORT QUEUE *mgpuStreamCreate() {
+
+ return catchAll([&]() { return new QUEUE(); });
+}
+
+extern "C" SYCL_RUNTIME_EXPORT void mgpuStreamDestroy(QUEUE *queue) {
+ catchAll([&]() { delete queue; });
+}
+
+extern "C" SYCL_RUNTIME_EXPORT void *mgpuMemAlloc(uint64_t size, QUEUE *queue,
+ bool isShared) {
+ return catchAll([&]() {
+ return allocDeviceMemory(queue, static_cast<size_t>(size), true);
+ });
+}
+
+extern "C" SYCL_RUNTIME_EXPORT void mgpuMemFree(void *ptr, QUEUE *queue) {
+ catchAll([&]() {
+ if (ptr) {
+ deallocDeviceMemory(queue, ptr);
+ }
+ });
+}
+
+extern "C" SYCL_RUNTIME_EXPORT ze_module_handle_t
+mgpuModuleLoad(const void *data, size_t gpuBlobSize) {
+ return catchAll([&]() { return loadModule(data, gpuBlobSize); });
+}
+
+extern "C" SYCL_RUNTIME_EXPORT sycl::kernel *
+mgpuModuleGetFunction(ze_module_handle_t module, const char *name) {
+ return catchAll([&]() { return getKernel(module, name); });
+}
+
+extern "C" SYCL_RUNTIME_EXPORT void
+mgpuLaunchKernel(sycl::kernel *kernel, size_t gridX, size_t gridY, size_t gridZ,
+ size_t blockX, size_t blockY, size_t blockZ,
+ size_t sharedMemBytes, QUEUE *queue, void **params,
+ void **extra, size_t paramsCount) {
+ return catchAll([&]() {
+ launchKernel(queue, kernel, gridX, gridY, gridZ, blockX, blockY, blockZ,
+ sharedMemBytes, params, paramsCount);
+ });
+}
+
+extern "C" SYCL_RUNTIME_EXPORT void mgpuStreamSynchronize(QUEUE *queue) {
+
+ catchAll([&]() { queue->syclQueue_.wait(); });
+}
+
+extern "C" SYCL_RUNTIME_EXPORT void
+mgpuModuleUnload(ze_module_handle_t module) {
+
+ catchAll([&]() { L0_SAFE_CALL(zeModuleDestroy(module)); });
+}
\ No newline at end of file
>From 4549a9ff41c271c66af89cf74ff26de8e87ae1b6 Mon Sep 17 00:00:00 2001
From: Nishant Patel <nishant.b.patel at intel.com>
Date: Thu, 19 Oct 2023 21:50:15 +0000
Subject: [PATCH 2/6] add source file to cmake
---
mlir/lib/ExecutionEngine/CMakeLists.txt | 1 +
1 file changed, 1 insertion(+)
diff --git a/mlir/lib/ExecutionEngine/CMakeLists.txt b/mlir/lib/ExecutionEngine/CMakeLists.txt
index ea33c2c6ed261e1..673efa368213983 100644
--- a/mlir/lib/ExecutionEngine/CMakeLists.txt
+++ b/mlir/lib/ExecutionEngine/CMakeLists.txt
@@ -12,6 +12,7 @@ set(LLVM_OPTIONAL_SOURCES
RunnerUtils.cpp
OptUtils.cpp
JitRunner.cpp
+ SyclRuntimeWrappers.cpp
)
# Use a separate library for OptUtils, to avoid pulling in the entire JIT and
>From 95b9d423a9e4fc0f317e2468696a9dafefd20f46 Mon Sep 17 00:00:00 2001
From: "Lee, Sang Ik" <sang.ik.lee at intel.com>
Date: Thu, 19 Oct 2023 15:01:26 -0700
Subject: [PATCH 3/6] Add build support.
---
mlir/CMakeLists.txt | 1 +
mlir/cmake/modules/FindLevelZero.cmake | 221 +++++++++++++++++++++++
mlir/cmake/modules/FindSyclRuntime.cmake | 68 +++++++
mlir/lib/ExecutionEngine/CMakeLists.txt | 35 ++++
4 files changed, 325 insertions(+)
create mode 100644 mlir/cmake/modules/FindLevelZero.cmake
create mode 100644 mlir/cmake/modules/FindSyclRuntime.cmake
diff --git a/mlir/CMakeLists.txt b/mlir/CMakeLists.txt
index ac120aad0d1eda7..16ff950089734b7 100644
--- a/mlir/CMakeLists.txt
+++ b/mlir/CMakeLists.txt
@@ -126,6 +126,7 @@ add_definitions(-DMLIR_ROCM_CONVERSIONS_ENABLED=${MLIR_ENABLE_ROCM_CONVERSIONS})
set(MLIR_ENABLE_DEPRECATED_GPU_SERIALIZATION 0 CACHE BOOL "Enable deprecated GPU serialization passes")
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")
set(MLIR_ENABLE_SPIRV_CPU_RUNNER 0 CACHE BOOL "Enable building the mlir SPIR-V cpu runner")
set(MLIR_ENABLE_VULKAN_RUNNER 0 CACHE BOOL "Enable building the mlir Vulkan runner")
set(MLIR_ENABLE_NVPTXCOMPILER 0 CACHE BOOL
diff --git a/mlir/cmake/modules/FindLevelZero.cmake b/mlir/cmake/modules/FindLevelZero.cmake
new file mode 100644
index 000000000000000..012187f0afc0b07
--- /dev/null
+++ b/mlir/cmake/modules/FindLevelZero.cmake
@@ -0,0 +1,221 @@
+# CMake find_package() module for level-zero
+#
+# Example usage:
+#
+# find_package(LevelZero)
+#
+# If successful, the following variables will be defined:
+# LevelZero_FOUND
+# LevelZero_INCLUDE_DIRS
+# LevelZero_LIBRARY
+# LevelZero_LIBRARIES_DIR
+#
+# By default, the module searches the standard paths to locate the "ze_api.h"
+# and the ze_loader shared library. When using a custom level-zero installation,
+# the environment variable "LEVEL_ZERO_DIR" should be specified telling the
+# module to get the level-zero library and headers from that location.
+
+include(FindPackageHandleStandardArgs)
+
+# Search path priority
+# 1. CMake Variable LEVEL_ZERO_DIR
+# 2. Environment Variable LEVEL_ZERO_DIR
+
+if(NOT LEVEL_ZERO_DIR)
+ if(DEFINED ENV{LEVEL_ZERO_DIR})
+ set(LEVEL_ZERO_DIR "$ENV{LEVEL_ZERO_DIR}")
+ endif()
+endif()
+
+if(LEVEL_ZERO_DIR)
+ find_path(LevelZero_INCLUDE_DIR
+ NAMES level_zero/ze_api.h
+ PATHS ${LEVEL_ZERO_DIR}/include
+ NO_DEFAULT_PATH
+ )
+
+ if(LINUX)
+ find_library(LevelZero_LIBRARY
+ NAMES ze_loader
+ PATHS ${LEVEL_ZERO_DIR}/lib
+ ${LEVEL_ZERO_DIR}/lib/x86_64-linux-gnu
+ NO_DEFAULT_PATH
+ )
+ else()
+ find_library(LevelZero_LIBRARY
+ NAMES ze_loader
+ PATHS ${LEVEL_ZERO_DIR}/lib
+ NO_DEFAULT_PATH
+ )
+ endif()
+else()
+ find_path(LevelZero_INCLUDE_DIR
+ NAMES level_zero/ze_api.h
+ )
+
+ find_library(LevelZero_LIBRARY
+ NAMES ze_loader
+ )
+endif()
+
+# Compares the two version string that are supposed to be in x.y.z format
+# and reports if the argument VERSION_STR1 is greater than or equal than
+# version_str2. The strings are compared lexicographically after conversion to
+# lists of equal lengths, with the shorter string getting zero-padded.
+function(compare_versions VERSION_STR1 VERSION_STR2 OUTPUT)
+ # Convert the strings to list
+ string(REPLACE "." ";" VL1 ${VERSION_STR1})
+ string(REPLACE "." ";" VL2 ${VERSION_STR2})
+ # get lengths of both lists
+ list(LENGTH VL1 VL1_LEN)
+ list(LENGTH VL2 VL2_LEN)
+ set(LEN ${VL1_LEN})
+ # If they differ in size pad the shorter list with 0s
+ if(VL1_LEN GREATER VL2_LEN)
+ math(EXPR DIFF "${VL1_LEN} - ${VL2_LEN}" OUTPUT_FORMAT DECIMAL)
+ foreach(IDX RANGE 1 ${DIFF} 1)
+ list(APPEND VL2 "0")
+ endforeach()
+ elseif(VL2_LEN GREATER VL2_LEN)
+ math(EXPR DIFF "${VL1_LEN} - ${VL2_LEN}" OUTPUT_FORMAT DECIMAL)
+ foreach(IDX RANGE 1 ${DIFF} 1)
+ list(APPEND VL2 "0")
+ endforeach()
+ set(LEN ${VL2_LEN})
+ endif()
+ math(EXPR LEN_SUB_ONE "${LEN}-1")
+ foreach(IDX RANGE 0 ${LEN_SUB_ONE} 1)
+ list(GET VL1 ${IDX} VAL1)
+ list(GET VL2 ${IDX} VAL2)
+
+ if(${VAL1} GREATER ${VAL2})
+ set(${OUTPUT} TRUE PARENT_SCOPE)
+ break()
+ elseif(${VAL1} LESS ${VAL2})
+ set(${OUTPUT} FALSE PARENT_SCOPE)
+ break()
+ else()
+ set(${OUTPUT} TRUE PARENT_SCOPE)
+ endif()
+ endforeach()
+
+ endfunction(compare_versions)
+
+# Creates a small function to run and extract the LevelZero loader version.
+function(get_l0_loader_version)
+
+ set(L0_VERSIONEER_SRC
+ [====[
+ #include <iostream>
+ #include <level_zero/loader/ze_loader.h>
+ #include <string>
+ int main() {
+ ze_result_t result;
+ std::string loader("loader");
+ zel_component_version_t *versions;
+ size_t size = 0;
+ result = zeInit(0);
+ if (result != ZE_RESULT_SUCCESS) {
+ std::cerr << "Failed to init ze driver" << std::endl;
+ return -1;
+ }
+ zelLoaderGetVersions(&size, nullptr);
+ versions = new zel_component_version_t[size];
+ zelLoaderGetVersions(&size, versions);
+ for (size_t i = 0; i < size; i++) {
+ if (loader.compare(versions[i].component_name) == 0) {
+ std::cout << versions[i].component_lib_version.major << "."
+ << versions[i].component_lib_version.minor << "."
+ << versions[i].component_lib_version.patch;
+ break;
+ }
+ }
+ delete[] versions;
+ return 0;
+ }
+ ]====]
+ )
+
+ set(L0_VERSIONEER_FILE ${CMAKE_BINARY_DIR}/temp/l0_versioneer.cpp)
+
+ file(WRITE ${L0_VERSIONEER_FILE} "${L0_VERSIONEER_SRC}")
+
+ # We need both the directories in the include path as ze_loader.h
+ # includes "ze_api.h" and not "level_zero/ze_api.h".
+ list(APPEND INCLUDE_DIRS ${LevelZero_INCLUDE_DIR})
+ list(APPEND INCLUDE_DIRS ${LevelZero_INCLUDE_DIR}/level_zero)
+ list(JOIN INCLUDE_DIRS ";" INCLUDE_DIRS_STR)
+ try_run(L0_VERSIONEER_RUN L0_VERSIONEER_COMPILE
+ "${CMAKE_BINARY_DIR}"
+ "${L0_VERSIONEER_FILE}"
+ LINK_LIBRARIES ${LevelZero_LIBRARY}
+ CMAKE_FLAGS
+ "-DINCLUDE_DIRECTORIES=${INCLUDE_DIRS_STR}"
+ RUN_OUTPUT_VARIABLE L0_VERSION
+ )
+ if(${L0_VERSIONEER_COMPILE} AND (DEFINED L0_VERSIONEER_RUN))
+ set(LevelZero_VERSION ${L0_VERSION} PARENT_SCOPE)
+ message(STATUS "Found Level Zero of version: ${L0_VERSION}")
+ else()
+ message(FATAL_ERROR
+ "Could not compile a level-zero program to extract loader version"
+ )
+ endif()
+endfunction(get_l0_loader_version)
+
+if(LevelZero_INCLUDE_DIR AND LevelZero_LIBRARY)
+ list(APPEND LevelZero_LIBRARIES "${LevelZero_LIBRARY}")
+ list(APPEND LevelZero_INCLUDE_DIRS ${LevelZero_INCLUDE_DIR})
+ if(OpenCL_FOUND)
+ list(APPEND LevelZero_INCLUDE_DIRS ${OpenCL_INCLUDE_DIRS})
+ endif()
+
+ cmake_path(GET LevelZero_LIBRARY PARENT_PATH LevelZero_LIBRARIES_PATH)
+ set(LevelZero_LIBRARIES_DIR ${LevelZero_LIBRARIES_PATH})
+
+ if(NOT TARGET LevelZero::LevelZero)
+ add_library(LevelZero::LevelZero INTERFACE IMPORTED)
+ set_target_properties(LevelZero::LevelZero
+ PROPERTIES INTERFACE_LINK_LIBRARIES "${LevelZero_LIBRARIES}"
+ )
+ set_target_properties(LevelZero::LevelZero
+ PROPERTIES INTERFACE_INCLUDE_DIRECTORIES "${LevelZero_INCLUDE_DIRS}"
+ )
+ endif()
+endif()
+
+# Check if a specific version of Level Zero is required
+if(LevelZero_FIND_VERSION)
+ get_l0_loader_version()
+ set(VERSION_GT_FIND_VERSION FALSE)
+ compare_versions(
+ ${LevelZero_VERSION}
+ ${LevelZero_FIND_VERSION}
+ VERSION_GT_FIND_VERSION
+ )
+ if(${VERSION_GT_FIND_VERSION})
+ set(LevelZero_FOUND TRUE)
+ else()
+ set(LevelZero_FOUND FALSE)
+ endif()
+else()
+ set(LevelZero_FOUND TRUE)
+endif()
+
+find_package_handle_standard_args(LevelZero
+ REQUIRED_VARS
+ LevelZero_FOUND
+ LevelZero_INCLUDE_DIRS
+ LevelZero_LIBRARY
+ LevelZero_LIBRARIES_DIR
+ HANDLE_COMPONENTS
+)
+mark_as_advanced(LevelZero_LIBRARY LevelZero_INCLUDE_DIRS)
+
+if(LevelZero_FOUND)
+ find_package_message(LevelZero "Found LevelZero: ${LevelZero_LIBRARY}"
+ "(found version ${LevelZero_VERSION})"
+ )
+else()
+ find_package_message(LevelZero "Could not find LevelZero" "")
+endif()
diff --git a/mlir/cmake/modules/FindSyclRuntime.cmake b/mlir/cmake/modules/FindSyclRuntime.cmake
new file mode 100644
index 000000000000000..38b065a3f284c2c
--- /dev/null
+++ b/mlir/cmake/modules/FindSyclRuntime.cmake
@@ -0,0 +1,68 @@
+# CMake find_package() module for SYCL Runtime
+#
+# Example usage:
+#
+# find_package(SyclRuntime)
+#
+# If successful, the following variables will be defined:
+# SyclRuntime_FOUND
+# SyclRuntime_INCLUDE_DIRS
+# SyclRuntime_LIBRARY
+# SyclRuntime_LIBRARIES_DIR
+#
+
+include(FindPackageHandleStandardArgs)
+
+if(NOT DEFINED ENV{CMPLR_ROOT})
+ message(WARNING "Please make sure to install Intel DPC++ Compiler and run setvars.(sh/bat)")
+ message(WARNING "You can download standalone Intel DPC++ Compiler from https://www.intel.com/content/www/us/en/developer/articles/tool/oneapi-standalone-components.html#compilers")
+else()
+ if(LINUX OR (${CMAKE_SYSTEM_NAME} MATCHES "Linux"))
+ set(SyclRuntime_ROOT "$ENV{CMPLR_ROOT}/linux")
+ elseif(WIN32)
+ set(SyclRuntime_ROOT "$ENV{CMPLR_ROOT}/windows")
+ endif()
+ list(APPEND SyclRuntime_INCLUDE_DIRS "${SyclRuntime_ROOT}/include")
+ list(APPEND SyclRuntime_INCLUDE_DIRS "${SyclRuntime_ROOT}/include/sycl")
+
+ set(SyclRuntime_LIBRARY_DIR "${SyclRuntime_ROOT}/lib")
+
+ message(STATUS "SyclRuntime_LIBRARY_DIR: ${SyclRuntime_LIBRARY_DIR}")
+ find_library(SyclRuntime_LIBRARY
+ NAMES sycl
+ PATHS ${SyclRuntime_LIBRARY_DIR}
+ NO_DEFAULT_PATH
+ )
+endif()
+
+if(SyclRuntime_LIBRARY)
+ set(SyclRuntime_FOUND TRUE)
+ if(NOT TARGET SyclRuntime::SyclRuntime)
+ add_library(SyclRuntime::SyclRuntime INTERFACE IMPORTED)
+ set_target_properties(SyclRuntime::SyclRuntime
+ PROPERTIES INTERFACE_LINK_LIBRARIES "${SyclRuntime_LIBRARY}"
+ )
+ set_target_properties(SyclRuntime::SyclRuntime
+ PROPERTIES INTERFACE_INCLUDE_DIRECTORIES "${SyclRuntime_INCLUDE_DIRS}"
+ )
+ endif()
+else()
+ set(SyclRuntime_FOUND FALSE)
+endif()
+
+find_package_handle_standard_args(SyclRuntime
+ REQUIRED_VARS
+ SyclRuntime_FOUND
+ SyclRuntime_INCLUDE_DIRS
+ SyclRuntime_LIBRARY
+ SyclRuntime_LIBRARY_DIR
+ HANDLE_COMPONENTS
+)
+
+mark_as_advanced(SyclRuntime_LIBRARY SyclRuntime_INCLUDE_DIRS)
+
+if(SyclRuntime_FOUND)
+ find_package_message(SyclRuntime "Found SyclRuntime: ${SyclRuntime_LIBRARY}" "")
+else()
+ find_package_message(SyclRuntime "Could not find SyclRuntime" "")
+endif()
diff --git a/mlir/lib/ExecutionEngine/CMakeLists.txt b/mlir/lib/ExecutionEngine/CMakeLists.txt
index 673efa368213983..fdc797763ae3a41 100644
--- a/mlir/lib/ExecutionEngine/CMakeLists.txt
+++ b/mlir/lib/ExecutionEngine/CMakeLists.txt
@@ -329,4 +329,39 @@ if(LLVM_ENABLE_PIC)
hip::host hip::amdhip64
)
endif()
+
+ if(MLIR_ENABLE_SYCL_RUNNER)
+ find_package(SyclRuntime)
+
+ if(NOT SyclRuntime_FOUND)
+ message(FATAL_ERROR "syclRuntime not found. Please set check oneapi installation and run setvars.sh.")
+ endif()
+
+ find_package(LevelZero)
+
+ if(NOT LevelZero_FOUND)
+ message(FATAL_ERROR "LevelZero not found. Please set LEVEL_ZERO_DIR.")
+ endif()
+
+ add_mlir_library(mlir_sycl_runtime
+ SHARED
+ SyclRuntimeWrappers.cpp
+
+ EXCLUDE_FROM_LIBMLIR
+ )
+
+ check_cxx_compiler_flag("-frtti" CXX_HAS_FRTTI_FLAG)
+ if(NOT CXX_HAS_FRTTI_FLAG)
+ message(FATAL_ERROR "CXX compiler does not accept flag -frtti")
+ endif()
+ target_compile_options (mlir_sycl_runtime PUBLIC -fexceptions -frtti)
+
+ target_include_directories(mlir_sycl_runtime PRIVATE
+ ${MLIR_INCLUDE_DIRS}
+ )
+
+ target_link_libraries(mlir_sycl_runtime PRIVATE LevelZero::LevelZero SyclRuntime::SyclRuntime)
+
+ set_property(TARGET mlir_sycl_runtime APPEND PROPERTY BUILD_RPATH "${LevelZero_LIBRARIES_DIR}" "${SyclRuntime_LIBRARIES_DIR}")
+ endif()
endif()
>From 7d54485d936d84a8089b680605a0e8e3d8d09a63 Mon Sep 17 00:00:00 2001
From: Nishant Patel <nishant.b.patel at intel.com>
Date: Tue, 24 Oct 2023 20:24:05 +0000
Subject: [PATCH 4/6] Address PR feedback
---
mlir/lib/ExecutionEngine/SyclRuntimeWrappers.cpp | 6 +-----
1 file changed, 1 insertion(+), 5 deletions(-)
diff --git a/mlir/lib/ExecutionEngine/SyclRuntimeWrappers.cpp b/mlir/lib/ExecutionEngine/SyclRuntimeWrappers.cpp
index f3a082922d61485..fe947657ebc3b1e 100644
--- a/mlir/lib/ExecutionEngine/SyclRuntimeWrappers.cpp
+++ b/mlir/lib/ExecutionEngine/SyclRuntimeWrappers.cpp
@@ -6,13 +6,12 @@
//
//===----------------------------------------------------------------------===//
//
-// Implements C wrappers around the sycl runtime library.
+// Implements wrappers around the sycl runtime library with C linkage
//
//===----------------------------------------------------------------------===//
#include <algorithm>
#include <array>
-#include <atomic>
#include <cassert>
#include <cfloat>
#include <cstdint>
@@ -20,12 +19,9 @@
#include <cstdlib>
#include <stdexcept>
#include <tuple>
-#include <vector>
#include <CL/sycl.hpp>
#include <level_zero/ze_api.h>
-#include <map>
-#include <mutex>
#include <sycl/ext/oneapi/backend/level_zero.hpp>
#ifdef _WIN32
>From e9fa1ca9af0688ceacadfcc1045e2ac75dfd64b4 Mon Sep 17 00:00:00 2001
From: Nishant Patel <nishant.b.patel at intel.com>
Date: Wed, 25 Oct 2023 22:22:18 +0000
Subject: [PATCH 5/6] clang-format style
---
mlir/lib/ExecutionEngine/SyclRuntimeWrappers.cpp | 3 +--
1 file changed, 1 insertion(+), 2 deletions(-)
diff --git a/mlir/lib/ExecutionEngine/SyclRuntimeWrappers.cpp b/mlir/lib/ExecutionEngine/SyclRuntimeWrappers.cpp
index c36ae2859ea65ab..6912e2c5586bca2 100644
--- a/mlir/lib/ExecutionEngine/SyclRuntimeWrappers.cpp
+++ b/mlir/lib/ExecutionEngine/SyclRuntimeWrappers.cpp
@@ -10,7 +10,6 @@
//
//===----------------------------------------------------------------------===//
-
#include <CL/sycl.hpp>
#include <level_zero/ze_api.h>
#include <sycl/ext/oneapi/backend/level_zero.hpp>
@@ -71,7 +70,7 @@ static sycl::device getDefaultDevice() {
}
static sycl::context getDefaultContext() {
- static sycl::context syclContext {getDefaultDevice()};
+ static sycl::context syclContext{getDefaultDevice()};
return syclContext;
}
>From 1af4b627b5ac642bac61984cce118a77a5cf6bbf Mon Sep 17 00:00:00 2001
From: Nishant Patel <nishant.b.patel at intel.com>
Date: Thu, 26 Oct 2023 15:56:04 +0000
Subject: [PATCH 6/6] Comment out unused argument
---
mlir/lib/ExecutionEngine/SyclRuntimeWrappers.cpp | 2 +-
1 file changed, 1 insertion(+), 1 deletion(-)
diff --git a/mlir/lib/ExecutionEngine/SyclRuntimeWrappers.cpp b/mlir/lib/ExecutionEngine/SyclRuntimeWrappers.cpp
index 6912e2c5586bca2..c250340c38fc77d 100644
--- a/mlir/lib/ExecutionEngine/SyclRuntimeWrappers.cpp
+++ b/mlir/lib/ExecutionEngine/SyclRuntimeWrappers.cpp
@@ -190,7 +190,7 @@ extern "C" SYCL_RUNTIME_EXPORT void
mgpuLaunchKernel(sycl::kernel *kernel, size_t gridX, size_t gridY, size_t gridZ,
size_t blockX, size_t blockY, size_t blockZ,
size_t sharedMemBytes, sycl::queue *queue, void **params,
- void **extra, size_t paramsCount) {
+ void ** /*extra*/, size_t paramsCount) {
return catchAll([&]() {
launchKernel(queue, kernel, gridX, gridY, gridZ, blockX, blockY, blockZ,
sharedMemBytes, params, paramsCount);
More information about the Mlir-commits
mailing list