From 5ccd81d920fefe92bb3d5dde12cbd0a65aad2034 Mon Sep 17 00:00:00 2001 From: Alberto Cabrera Date: Thu, 14 Mar 2024 14:12:38 +0000 Subject: [PATCH 1/8] [SYCL][COMPAT] Added Check Error Macro. Updated macro definitions. --- sycl/doc/syclcompat/README.md | 39 +++++++++++++++++++++++---- sycl/include/syclcompat/defs.hpp | 44 ++++++++++++++++++++++++++----- sycl/test-e2e/syclcompat/defs.cpp | 33 ++++++++++++++++++++--- 3 files changed, 101 insertions(+), 15 deletions(-) diff --git a/sycl/doc/syclcompat/README.md b/sycl/doc/syclcompat/README.md index 9b72013f6ff86..d1f0fddc9ee46 100644 --- a/sycl/doc/syclcompat/README.md +++ b/sycl/doc/syclcompat/README.md @@ -1177,16 +1177,45 @@ kernel names during machine translation. `get_sycl_language_version` returns an integer representing the version of the SYCL spec supported by the current SYCL compiler. +The `SYCLCOMPAT_CHECK_ERROR` macro encapsulates an error-handling mechanism for +expressions that might throw `sycl::exception`. If no exceptions are thrown, it +returns `syclcompat::error_code::SUCCESS`. If an exception is caught, it prints +the error message to the standard error stream and returns +`syclcompat::error_code::DEFAULT_ERROR`. + ``` c++ namespace syclcompat { -#define __sycl_compat_align__(n) __attribute__((aligned(n))) -#define __sycl_compat_inline__ __inline__ __attribute__((always_inline)) +template class syclcompat_kernel_name; +template class syclcompat_kernel_scalar; + +#if defined(_MSC_VER) +#define __syclcompat_align__(n) __declspec(align(n)) +#define __syclcompat_inline__ __forceinline +#else +#define __syclcompat_align__(n) __attribute__((aligned(n))) +#define __syclcompat_inline__ __inline__ __attribute__((always_inline)) +#endif + +#if defined(_MSC_VER) +#define __syclcompat_noinline__ __declspec(noinline) +#else +#define __syclcompat_noinline__ __attribute__((noinline)) +#endif + +#define SYCLCOMPAT_COMPATIBILITY_TEMP (600) -#define __sycl_compat_noinline__ __attribute__((noinline)) +#ifdef _WIN32 +#define SYCLCOMPAT_EXPORT __declspec(dllexport) +#else +#define SYCLCOMPAT_EXPORT +#endif + +namespace syclcompat { +enum error_code { SUCCESS = 0, DEFAULT_ERROR = 999 }; +} -template class sycl_compat_kernel_name; -template class sycl_compat_kernel_scalar; +#define SYCLCOMPAT_CHECK_ERROR(expr) int get_sycl_language_version(); diff --git a/sycl/include/syclcompat/defs.hpp b/sycl/include/syclcompat/defs.hpp index 6e4d76cf99403..d2922bf441f2e 100644 --- a/sycl/include/syclcompat/defs.hpp +++ b/sycl/include/syclcompat/defs.hpp @@ -32,12 +32,44 @@ #pragma once -template class sycl_compat_kernel_name; -template class sycl_compat_kernel_scalar; +#include -#define __sycl_compat_align__(n) alignas(n) -#define __sycl_compat_inline__ __inline__ __attribute__((always_inline)) +template class syclcompat_kernel_name; +template class syclcompat_kernel_scalar; -#define __sycl_compat_noinline__ __attribute__((noinline)) +#if defined(_MSC_VER) +#define __syclcompat_align__(n) __declspec(align(n)) +#define __syclcompat_inline__ __forceinline +#else +#define __syclcompat_align__(n) __attribute__((aligned(n))) +#define __syclcompat_inline__ __inline__ __attribute__((always_inline)) +#endif -#define SYCL_COMPAT_COMPATIBILITY_TEMP (600) +#if defined(_MSC_VER) +#define __syclcompat_noinline__ __declspec(noinline) +#else +#define __syclcompat_noinline__ __attribute__((noinline)) +#endif + +#define SYCLCOMPAT_COMPATIBILITY_TEMP (600) + +#ifdef _WIN32 +#define SYCLCOMPAT_EXPORT __declspec(dllexport) +#else +#define SYCLCOMPAT_EXPORT +#endif + +namespace syclcompat { +enum error_code { SUCCESS = 0, DEFAULT_ERROR = 999 }; +} + +#define SYCLCOMPAT_CHECK_ERROR(expr) \ + [&]() { \ + try { \ + expr; \ + return syclcompat::error_code::SUCCESS; \ + } catch (sycl::exception const &e) { \ + std::cerr << e.what() << std::endl; \ + return syclcompat::error_code::DEFAULT_ERROR; \ + } \ + }() diff --git a/sycl/test-e2e/syclcompat/defs.cpp b/sycl/test-e2e/syclcompat/defs.cpp index 01681d13610f7..dc34cf4f298c2 100644 --- a/sycl/test-e2e/syclcompat/defs.cpp +++ b/sycl/test-e2e/syclcompat/defs.cpp @@ -17,22 +17,47 @@ * Defs.cpp * * Description: - * __sycl_compat_align__ tests + * Syclcompat macros tests **************************************************************************/ // RUN: %clangxx -fsycl %s -o %t.out // RUN: %{run} %t.out #include +#include + +#include + #include -int main() { - struct __sycl_compat_align__(16) { +void test_align() { + std::cout << __PRETTY_FUNCTION__ << std::endl; + + constexpr std::size_t expected_size = 16; + struct __syclcompat_align__(expected_size) { int a; char c; } s; - assert(sizeof(s) == 16); + assert(sizeof(s) == expected_size); +} + +void test_check_error() { + std::cout << __PRETTY_FUNCTION__ << std::endl; + + auto error_throw = []() { + throw sycl::exception(sycl::make_error_code(sycl::errc::invalid), + "Expected invalid exception in test_check_error"); + }; + + assert(syclcompat::error_code::SUCCESS == SYCLCOMPAT_CHECK_ERROR()); + assert(syclcompat::error_code::DEFAULT_ERROR == + SYCLCOMPAT_CHECK_ERROR(error_throw())); +} + +int main() { + test_align(); + test_check_error(); return 0; } From 23f90fc68c4434ecd6ea5230c303b516eb85a1b4 Mon Sep 17 00:00:00 2001 From: Alberto Cabrera Date: Thu, 14 Mar 2024 14:17:40 +0000 Subject: [PATCH 2/8] [SYCL][COMPAT] Reduced verbosity of macro definitions --- sycl/include/syclcompat/defs.hpp | 6 +----- 1 file changed, 1 insertion(+), 5 deletions(-) diff --git a/sycl/include/syclcompat/defs.hpp b/sycl/include/syclcompat/defs.hpp index d2922bf441f2e..bb93b24920ad3 100644 --- a/sycl/include/syclcompat/defs.hpp +++ b/sycl/include/syclcompat/defs.hpp @@ -40,14 +40,10 @@ template class syclcompat_kernel_scalar; #if defined(_MSC_VER) #define __syclcompat_align__(n) __declspec(align(n)) #define __syclcompat_inline__ __forceinline +#define __syclcompat_noinline__ __declspec(noinline) #else #define __syclcompat_align__(n) __attribute__((aligned(n))) #define __syclcompat_inline__ __inline__ __attribute__((always_inline)) -#endif - -#if defined(_MSC_VER) -#define __syclcompat_noinline__ __declspec(noinline) -#else #define __syclcompat_noinline__ __attribute__((noinline)) #endif From 081ead19e1fa93491aff936ebb996b525750a918 Mon Sep 17 00:00:00 2001 From: Alberto Cabrera Date: Mon, 18 Mar 2024 11:26:06 +0000 Subject: [PATCH 3/8] [SYCL][COMPAT] SYCLCOMPAT_CHECK_ERROR split for backend and library errors --- sycl/doc/syclcompat/README.md | 11 +++++++---- sycl/include/syclcompat/defs.hpp | 5 ++++- sycl/test-e2e/syclcompat/defs.cpp | 10 ++++++++-- 3 files changed, 19 insertions(+), 7 deletions(-) diff --git a/sycl/doc/syclcompat/README.md b/sycl/doc/syclcompat/README.md index d1f0fddc9ee46..b7c82145efa38 100644 --- a/sycl/doc/syclcompat/README.md +++ b/sycl/doc/syclcompat/README.md @@ -1178,10 +1178,13 @@ kernel names during machine translation. SYCL spec supported by the current SYCL compiler. The `SYCLCOMPAT_CHECK_ERROR` macro encapsulates an error-handling mechanism for -expressions that might throw `sycl::exception`. If no exceptions are thrown, it -returns `syclcompat::error_code::SUCCESS`. If an exception is caught, it prints -the error message to the standard error stream and returns -`syclcompat::error_code::DEFAULT_ERROR`. +expressions that might throw `sycl::exception` and `std::runtime_error`. +If no exceptions are thrown, it returns `syclcompat::error_code::SUCCESS`. +If a `sycl::exception` is caught, it returns `syclcompat::error_code::BACKEND_ERROR`. +If a `std::runtime_error` exception is caught, +`syclcompat::error_code::DEFAULT_ERROR` is returned instead. +For both cases, it prints the error message to the standar +d error stream. ``` c++ namespace syclcompat { diff --git a/sycl/include/syclcompat/defs.hpp b/sycl/include/syclcompat/defs.hpp index bb93b24920ad3..703027f66bd60 100644 --- a/sycl/include/syclcompat/defs.hpp +++ b/sycl/include/syclcompat/defs.hpp @@ -56,7 +56,7 @@ template class syclcompat_kernel_scalar; #endif namespace syclcompat { -enum error_code { SUCCESS = 0, DEFAULT_ERROR = 999 }; +enum error_code { SUCCESS = 0, BACKEND_ERROR = 1, DEFAULT_ERROR = 999 }; } #define SYCLCOMPAT_CHECK_ERROR(expr) \ @@ -65,6 +65,9 @@ enum error_code { SUCCESS = 0, DEFAULT_ERROR = 999 }; expr; \ return syclcompat::error_code::SUCCESS; \ } catch (sycl::exception const &e) { \ + std::cerr << e.what() << std::endl; \ + return syclcompat::error_code::BACKEND_ERROR; \ + } catch (std::runtime_error const &e) { \ std::cerr << e.what() << std::endl; \ return syclcompat::error_code::DEFAULT_ERROR; \ } \ diff --git a/sycl/test-e2e/syclcompat/defs.cpp b/sycl/test-e2e/syclcompat/defs.cpp index dc34cf4f298c2..19dd97347f051 100644 --- a/sycl/test-e2e/syclcompat/defs.cpp +++ b/sycl/test-e2e/syclcompat/defs.cpp @@ -45,14 +45,20 @@ void test_align() { void test_check_error() { std::cout << __PRETTY_FUNCTION__ << std::endl; - auto error_throw = []() { + auto sycl_error_throw = []() { throw sycl::exception(sycl::make_error_code(sycl::errc::invalid), "Expected invalid exception in test_check_error"); }; + auto runtime_error_throw = []() { + throw std::runtime_error("Expected invalid exception in test_check_error"); + }; + assert(syclcompat::error_code::SUCCESS == SYCLCOMPAT_CHECK_ERROR()); + assert(syclcompat::error_code::BACKEND_ERROR == + SYCLCOMPAT_CHECK_ERROR(sycl_error_throw())); assert(syclcompat::error_code::DEFAULT_ERROR == - SYCLCOMPAT_CHECK_ERROR(error_throw())); + SYCLCOMPAT_CHECK_ERROR(runtime_error_throw())); } int main() { From 4ba37d22a544d9765909f5ad63617db45bb0bf19 Mon Sep 17 00:00:00 2001 From: Alberto Cabrera Date: Mon, 18 Mar 2024 11:19:10 +0000 Subject: [PATCH 4/8] [SYCL][COMPAT] Added library and function load helper functions --- sycl/doc/syclcompat/README.md | 37 +++ sycl/include/syclcompat/kernel.hpp | 412 ++++++++++++++++++++++++++++- 2 files changed, 446 insertions(+), 3 deletions(-) diff --git a/sycl/doc/syclcompat/README.md b/sycl/doc/syclcompat/README.md index b7c82145efa38..f4eca6d756518 100644 --- a/sycl/doc/syclcompat/README.md +++ b/sycl/doc/syclcompat/README.md @@ -1233,6 +1233,15 @@ to get the kernel information. Overloads are provided to allow either returning a `kernel_function_info` object, or to return by pointer argument. In the current version, `kernel_function_info` describes only maximum work-group size. +SYCLcompat also provides the `kernel_library` and `kernel_function` classes. +`kernel_library` facilitates the loading and unloading of kernel libraries. +`kernel_function` represents a specific kernel function within a loaded librariy +and can be invoked with specified arguments. +`load_kernel_library`, `load_kernel_library_mem`, and `unload_kernel_library` are +free functions to handle the loading and unloading of `kernel_library` objects. +`get_kernel_function`, and `invoke_kernel_function` offer a similar functionality +for `kernel_function` objects. + ``` c++ namespace syclcompat { @@ -1243,6 +1252,34 @@ struct kernel_function_info { static void get_kernel_function_info(kernel_function_info *kernel_info, const void *function); static kernel_function_info get_kernel_function_info(const void *function); + +class kernel_library { + kernel_library(); + kernel_library(void *ptr); + operator void *() const; +}; + +static kernel_library load_kernel_library(const std::string &name); +static kernel_library load_kernel_library_mem(char const *const image); +static void unload_kernel_library(const kernel_library &library); + +class kernel_function { + kernel_function(); + kernel_function(kernel_functor ptr); + operator void *() const; + void operator()(sycl::queue &q, const sycl::nd_range<3> &range, + unsigned int local_mem_size, void **args, void **extra); +}; + +static kernel_function get_kernel_function(kernel_library &library, + const std::string &name); +static void invoke_kernel_function(kernel_function &function, + sycl::queue &queue, + sycl::range<3> group_range, + sycl::range<3> local_range, + unsigned int local_mem_size, + void **kernelParams, void **extra); + } // namespace syclcompat ``` diff --git a/sycl/include/syclcompat/kernel.hpp b/sycl/include/syclcompat/kernel.hpp index 0132e98e43841..00a9808d1d1ba 100644 --- a/sycl/include/syclcompat/kernel.hpp +++ b/sycl/include/syclcompat/kernel.hpp @@ -31,6 +31,25 @@ #pragma once +#ifdef _WIN32 +#include +#include +#else +#include +#endif + +#if defined(__has_include) && __has_include() +#include +#elif defined(__has_include) && __has_include() +#include +#else +#error "SYCLomatic runtime requires C++ filesystem support" +#endif + +#include +#include +#include + #include #include #include @@ -44,14 +63,16 @@ struct kernel_function_info { int max_work_group_size = 0; }; -static void get_kernel_function_info(kernel_function_info *kernel_info, - const void *function) { +static inline void get_kernel_function_info(kernel_function_info *kernel_info, + const void *function) { kernel_info->max_work_group_size = detail::dev_mgr::instance() .current_device() .get_info(); } -static kernel_function_info get_kernel_function_info(const void *function) { + +static inline kernel_function_info +get_kernel_function_info(const void *function) { kernel_function_info kernel_info; kernel_info.max_work_group_size = detail::dev_mgr::instance() @@ -60,4 +81,389 @@ static kernel_function_info get_kernel_function_info(const void *function) { return kernel_info; } +namespace detail { + +#if defined(__has_include) && __has_include() +namespace fs = std::filesystem; +#else +namespace fs = std::experimental::filesystem; +#endif + +/// Write data to temporary file and return absolute path to temporary file. +/// Temporary file is created in a temporary directory both of which have random +/// names with only the user having access permissions. Only one temporary file +/// will be created in the temporary directory. +static inline fs::path write_data_to_file(char const *const data, size_t size) { + std::error_code ec; + + if (sizeof(size_t) >= sizeof(std::streamsize) && + size > (std::numeric_limits::max)()) + throw std::runtime_error("[SYCLcompat] data file too large"); + + // random number generator + std::random_device dev; + std::mt19937 prng(dev()); + std::uniform_int_distribution rand(0); + + // find temporary directory + auto tmp_dir = fs::temp_directory_path(ec); + if (ec) + throw std::runtime_error("[SYCLcompat] could not find temporary directory"); + + // create private directory + std::stringstream directory; + fs::path directory_path; + constexpr int max_attempts = 5; + int i; + + for (i = 0; i < max_attempts; i++) { + directory << std::hex << rand(prng); + directory_path = tmp_dir / directory.str(); + if (fs::create_directory(directory_path)) { + break; + } + } + if (i == max_attempts) + throw std::runtime_error("[SYCLcompat] could not create directory"); + + // only allow owner permissions to private directory + fs::permissions(directory_path, fs::perms::owner_all, ec); + if (ec) + throw std::runtime_error( + "[SYCLcompat] could not set directory permissions"); + + // random filename in private directory + std::stringstream filename; + filename << std::hex << rand(prng); +#ifdef _WIN32 + auto filepath = directory_path / (filename.str() + ".dll"); +#else + auto filepath = directory_path / filename.str(); +#endif + + // write data to temporary file + auto outfile = std::ofstream(filepath, std::ios::out | std::ios::binary); + if (outfile) { + // only allow program to write file + fs::permissions(filepath, fs::perms::owner_write, ec); + if (ec) + throw std::runtime_error("[SYCLcompat] could not set permissions"); + + outfile.write(data, size); + if (!outfile.good()) + throw std::runtime_error("[SYCLcompat] could not write data"); + outfile.close(); + + // only allow program to read/execute file + fs::permissions(filepath, fs::perms::owner_read | fs::perms::owner_exec, + ec); + if (ec) + throw std::runtime_error("[SYCLcompat] could not set permissions"); + } else + throw std::runtime_error("[SYCLcompat] could not write data"); + + // check temporary file contents + auto infile = std::ifstream(filepath, std::ios::in | std::ios::binary); + if (infile) { + bool mismatch = false; + size_t cnt = 0; + + while (1) { + char c; + infile.get(c); + if (infile.eof()) + break; + if (c != data[cnt++]) + mismatch = true; + } + if (cnt != size || mismatch) + throw std::runtime_error( + "[SYCLcompat] file contents not written correctly"); + } else + throw std::runtime_error("[SYCLcompat] could not validate file"); + + if (!filepath.is_absolute()) + throw std::runtime_error("[SYCLcompat] temporary filepath is not absolute"); + + return filepath; +} + +static inline uint16_t extract16(unsigned char const *const ptr) { + uint16_t ret = 0; + + ret |= static_cast(ptr[0]) << 0; + ret |= static_cast(ptr[1]) << 8; + + return (ret); +} + +static inline uint32_t extract32(unsigned char const *const ptr) { + uint32_t ret = 0; + + ret |= static_cast(ptr[0]) << 0; + ret |= static_cast(ptr[1]) << 8; + ret |= static_cast(ptr[2]) << 16; + ret |= static_cast(ptr[3]) << 24; + + return (ret); +} + +static inline uint64_t extract64(unsigned char const *const ptr) { + uint64_t ret = 0; + + ret |= static_cast(ptr[0]) << 0; + ret |= static_cast(ptr[1]) << 8; + ret |= static_cast(ptr[2]) << 16; + ret |= static_cast(ptr[3]) << 24; + ret |= static_cast(ptr[4]) << 32; + ret |= static_cast(ptr[5]) << 40; + ret |= static_cast(ptr[6]) << 48; + ret |= static_cast(ptr[7]) << 56; + + return (ret); +} + +static inline uint64_t get_lib_size(char const *const blob) { +#ifdef _WIN32 + /////////////////////////////////////////////////////////////////////// + // Analyze DOS stub + unsigned char const *const ublob = + reinterpret_cast(blob); + if (ublob[0] != 0x4d || ublob[1] != 0x5a) { + throw std::runtime_error("[SYCLcompat] blob is not a Windows DLL."); + } + uint32_t pe_header_offset = extract32(ublob + 0x3c); + + /////////////////////////////////////////////////////////////////////// + // Ananlyze PE-header + unsigned char const *const pe_header = ublob + pe_header_offset; + + // signature + uint32_t pe_signature = extract32(pe_header + 0); + if (pe_signature != 0x00004550) { + throw std::runtime_error( + "[SYCLcompat] PE-header signature is not 0x00004550"); + } + + // machine + uint16_t machine = extract16(pe_header + 4); + if (machine != 0x8664) { + throw std::runtime_error("[SYCLcompat] only DLLs for x64 supported"); + } + + // number of sections + uint16_t number_of_sections = extract16(pe_header + 6); + + // sizeof optional header + uint16_t sizeof_optional_header = extract16(pe_header + 20); + + // magic + uint16_t magic = extract16(pe_header + 24); + if (magic != 0x10b && magic != 0x20b) { + throw std::runtime_error("[SYCLcompat] MAGIC is not 0x010b or 0x020b"); + } + + /////////////////////////////////////////////////////////////////////// + // Analyze tail of optional header + constexpr int coff_header_size = 24; + + unsigned char const *const tail_of_optional_header = + pe_header + coff_header_size + sizeof_optional_header; + if (extract64(tail_of_optional_header - 8) != 0) { + throw std::runtime_error("Optional header not zero-padded"); + } + + /////////////////////////////////////////////////////////////////////// + // Analyze last section header + constexpr int section_header_size = 40; + unsigned char const *const last_section_header = + tail_of_optional_header + section_header_size * (number_of_sections - 1); + + uint32_t sizeof_raw_data = extract32(last_section_header + 16); + uint32_t pointer_to_raw_data = extract32(last_section_header + 20); + + return sizeof_raw_data + pointer_to_raw_data; +#else + if (blob[0] != 0x7F || blob[1] != 'E' || blob[2] != 'L' || blob[3] != 'F') + throw std::runtime_error("[SYCLcompat] blob is not in ELF format"); + + if (blob[4] != 0x02) + throw std::runtime_error("[SYCLcompat] only 64-bit headers are supported"); + + if (blob[5] != 0x01) + throw std::runtime_error( + "[SYCLcompat] only little-endian headers are supported"); + + unsigned char const *const ublob = + reinterpret_cast(blob); + uint64_t e_shoff = extract64(ublob + 0x28); + uint16_t e_shentsize = extract16(ublob + 0x3A); + uint16_t e_shnum = extract16(ublob + 0x3C); + + return e_shoff + (e_shentsize * e_shnum); +#endif +} + +#ifdef _WIN32 +class path_lib_record { +public: + void operator=(const path_lib_record &) = delete; + ~path_lib_record() { + for (auto entry : lib_to_path) { + FreeLibrary(static_cast(entry.first)); + fs::permissions(entry.second, fs::perms::owner_all); + fs::remove_all(entry.second.remove_filename()); + } + } + static void record_lib_path(fs::path path, void *library) { + lib_to_path[library] = path; + } + static void remove_lib(void *library) { + auto path = lib_to_path[library]; + std::error_code ec; + + FreeLibrary(static_cast(library)); + fs::permissions(path, fs::perms::owner_all); + if (fs::remove_all(path.remove_filename(), ec) != 2 || ec) + // one directory and one temporary file should have been deleted + throw std::runtime_error("[SYCLcompat] directory delete failed"); + + lib_to_path.erase(library); + } + +private: + static inline std::unordered_map lib_to_path; +}; +#endif + +} // namespace detail + +class kernel_library { +public: + kernel_library() : ptr{nullptr} {} + kernel_library(void *ptr) : ptr{ptr} {} + + operator void *() const { return ptr; } + +private: + void *ptr; +#ifdef _WIN32 + static inline detail::path_lib_record single_instance_to_trigger_destructor; +#endif +}; + +namespace detail { + +static inline kernel_library load_dl_from_data(char const *const data, + size_t size) { + fs::path filename = write_data_to_file(data, size); +#ifdef _WIN32 + void *so = LoadLibraryW(filename.wstring().c_str()); +#else + void *so = dlopen(filename.c_str(), RTLD_LAZY); +#endif + if (so == nullptr) + throw std::runtime_error("[SYCLcompat] failed to load kernel library"); + +#ifdef _WIN32 + detail::path_lib_record::record_lib_path(filename, so); +#else + std::error_code ec; + + // Windows DLL cannot be deleted while in use + if (fs::remove_all(filename.remove_filename(), ec) != 2 || ec) + // one directory and one temporary file should have been deleted + throw std::runtime_error("[SYCLcompat] directory delete failed"); +#endif + + return so; +} + +} // namespace detail + +/// Load kernel library and return a handle to use the library. +/// \param [in] name The name of the library. +static inline kernel_library load_kernel_library(const std::string &name) { + std::ifstream ifs; + ifs.open(name, std::ios::in | std::ios::binary); + + std::stringstream buffer; + buffer << ifs.rdbuf(); + + const std::string buffer_string = buffer.str(); + return detail::load_dl_from_data(buffer_string.c_str(), buffer_string.size()); +} + +/// Load kernel library whose image is alreay in memory and return a handle to +/// use the library. +/// \param [in] image A pointer to the image in memory. +static inline kernel_library load_kernel_library_mem(char const *const image) { + const size_t size = detail::get_lib_size(image); + + return detail::load_dl_from_data(image, size); +} + +/// Unload kernel library. +/// \param [in,out] library Handle to the library to be closed. +static inline void unload_kernel_library(const kernel_library &library) { +#ifdef _WIN32 + detail::path_lib_record::remove_lib(library); +#else + dlclose(library); +#endif +} + +class kernel_function { +public: + kernel_function() : ptr{nullptr} {} + kernel_function(kernel_functor ptr) : ptr{ptr} {} + + operator void *() const { return ((void *)ptr); } + + void operator()(sycl::queue &q, const sycl::nd_range<3> &range, + unsigned int local_mem_size, void **args, void **extra) { + ptr(q, range, local_mem_size, args, extra); + } + +private: + kernel_functor ptr; +}; + +/// Find kernel function in a kernel library and return its address. +/// \param [in] library Handle to the kernel library. +/// \param [in] name Name of the kernel function. +static inline kernel_function get_kernel_function(kernel_library &library, + const std::string &name) { +#ifdef _WIN32 + kernel_functor fn = reinterpret_cast( + GetProcAddress(static_cast(static_cast(library)), + (name + std::string("_wrapper")).c_str())); +#else + kernel_functor fn = reinterpret_cast( + dlsym(library, (name + std::string("_wrapper")).c_str())); +#endif + if (fn == nullptr) + throw std::runtime_error("[SYCLcompat] failed to get function"); + return fn; +} + +/// Invoke a kernel function. +/// \param [in] function kernel function. +/// \param [in] queue SYCL queue used to execute kernel +/// \param [in] group_range SYCL group range +/// \param [in] local_range SYCL local range +/// \param [in] local_mem_size The size of local memory required by the kernel +/// function. +/// \param [in] kernel_params Array of pointers to kernel arguments. +/// \param [in] extra Extra arguments. +static inline void invoke_kernel_function(kernel_function &function, + sycl::queue &queue, + sycl::range<3> group_range, + sycl::range<3> local_range, + unsigned int local_mem_size, + void **kernelParams, void **extra) { + function(queue, sycl::nd_range<3>(group_range * local_range, local_range), + local_mem_size, kernel_params, extra); +} + } // namespace syclcompat From 2bcb956f70f55a144b99845409a45695215f902b Mon Sep 17 00:00:00 2001 From: Alberto Cabrera Date: Mon, 18 Mar 2024 11:19:35 +0000 Subject: [PATCH 5/8] [SYCL][COMPAT] Updated kernel tests to operate using the new helpers --- .../kernel/Inputs/kernel_function.cpp | 120 +++++++++++------- .../kernel/Inputs/kernel_module.cpp | 13 +- 2 files changed, 77 insertions(+), 56 deletions(-) diff --git a/sycl/test-e2e/syclcompat/kernel/Inputs/kernel_function.cpp b/sycl/test-e2e/syclcompat/kernel/Inputs/kernel_function.cpp index 409bb8603adba..d3b4502df0544 100644 --- a/sycl/test-e2e/syclcompat/kernel/Inputs/kernel_function.cpp +++ b/sycl/test-e2e/syclcompat/kernel/Inputs/kernel_function.cpp @@ -40,8 +40,10 @@ #include +#include #include #include +#include template void testTemplateKernel(T *data) {} @@ -62,76 +64,98 @@ int getFuncAttrs() { return threadPerBlock; } -#ifdef WIN32 +void test_get_func_attrs() { + std::cout << __PRETTY_FUNCTION__ << std::endl; + syclcompat::device_ext &dev_ct1 = syclcompat::get_current_device(); -#define DECLARE_MODULE_VAR(var) HMODULE var -#define LOAD_LIB(lib) LoadLibraryA(lib) -#define LOAD_FUNCTOR(module, name) GetProcAddress(module, name) -#define FREE_LIB(module) FreeLibrary(module) + int size = dev_ct1.get_info(); + assert(getTemplateFuncAttrs() == size); + assert(getFuncAttrs() == size); +} -#else // LINUX +void call_library_func(syclcompat::kernel_library kernel_lib) { + sycl::queue q_ct1 = syclcompat::get_default_queue(); -#define DECLARE_MODULE_VAR(var) void *var -#define LOAD_LIB(lib) dlopen(lib, RTLD_LAZY) -#define LOAD_FUNCTOR(module, name) dlsym(module, name) -#define FREE_LIB(module) dlclose(module) + std::string FunctionName = "foo"; + syclcompat::kernel_function func; + SYCLCOMPAT_CHECK_ERROR( + func = syclcompat::get_kernel_function(kernel_lib, FunctionName.c_str())); -#endif + if (func == nullptr) { + std::cout << "Could not load function pointer" << std::endl << std::flush; + syclcompat::unload_kernel_library(kernel_lib); + assert(false); // FAIL + } -void test_kernel_functor_ptr() { - std::cout << __PRETTY_FUNCTION__ << std::endl; + int sharedSize = 10; + void **param = nullptr, **extra = nullptr; - syclcompat::device_ext &dev_ct1 = syclcompat::get_current_device(); - sycl::queue *q_ct1 = dev_ct1.default_queue(); + constexpr size_t NUM_ELEMENTS = 16; + int *dev = syclcompat::malloc(NUM_ELEMENTS); + syclcompat::fill(dev, 0, NUM_ELEMENTS); - int Size = dev_ct1.get_info(); - assert(getTemplateFuncAttrs() == Size); - assert(getFuncAttrs() == Size); + param = (void **)(&dev); + SYCLCOMPAT_CHECK_ERROR(syclcompat::invoke_kernel_function( + func, q_ct1, sycl::range<3>(1, 1, 2), sycl::range<3>(1, 1, 8), sharedSize, + param, extra)); + syclcompat::wait_and_throw(); + + int *host_mem = syclcompat::malloc_host(NUM_ELEMENTS); + syclcompat::memcpy(host_mem, dev, NUM_ELEMENTS); + for (int i = 0; i < NUM_ELEMENTS; i++) { + assert(host_mem[i] == i); + } - DECLARE_MODULE_VAR(M); - M = LOAD_LIB(TEST_SHARED_LIB); + SYCLCOMPAT_CHECK_ERROR(syclcompat::unload_kernel_library(kernel_lib)); - if (M == NULL) { + syclcompat::free(dev); + syclcompat::free(host_mem); +} + +void test_kernel_functor_ptr() { + std::cout << __PRETTY_FUNCTION__ << std::endl; + + syclcompat::kernel_library kernel_lib; + SYCLCOMPAT_CHECK_ERROR(kernel_lib = + syclcompat::load_kernel_library(TEST_SHARED_LIB)); + + if (kernel_lib == nullptr) { std::cout << "Could not load the library" << std::endl; std::cout << " " << TEST_SHARED_LIB << std::endl << std::flush; assert(false); // FAIL } - std::string FunctionName = "foo_wrapper"; - syclcompat::kernel_functor F; - F = (syclcompat::kernel_functor)LOAD_FUNCTOR(M, FunctionName.c_str()); + call_library_func(kernel_lib); +} - if (F == NULL) { - std::cout << "Could not load function pointer" << std::endl << std::flush; - FREE_LIB(M); - assert(false); // FAIL - } +void test_kernel_functor_ptr_memory() { + std::cout << __PRETTY_FUNCTION__ << std::endl; - int sharedSize = 10; - void **param = nullptr, **extra = nullptr; - if (!q_ct1->get_device().has(sycl::aspect::usm_shared_allocations)) - return; - int *dev = sycl::malloc_shared(16, *q_ct1); - for (int i = 0; i < 16; i++) { - dev[i] = 0; - } - param = (void **)(&dev); - F(*q_ct1, - sycl::nd_range<3>(sycl::range<3>(1, 1, 2) * sycl::range<3>(1, 1, 8), - sycl::range<3>(1, 1, 8)), - sharedSize, param, extra); - q_ct1->wait_and_throw(); - - for (int i = 0; i < 16; i++) { - assert(dev[i] == i); + sycl::queue q_ct1 = syclcompat::get_default_queue(); + + std::ifstream ifs; + ifs.open(TEST_SHARED_LIB, std::ios::in | std::ios::binary); + + std::stringstream buffer; + buffer << ifs.rdbuf(); + + syclcompat::kernel_library kernel_lib; + SYCLCOMPAT_CHECK_ERROR( + kernel_lib = syclcompat::load_kernel_library_mem(buffer.str().c_str())); + + if (kernel_lib == nullptr) { + std::cout << "Could not load the library" << std::endl; + std::cout << " " << TEST_SHARED_LIB << std::endl << std::flush; + assert(false); } - sycl::free(dev, *q_ct1); - FREE_LIB(M); + call_library_func(kernel_lib); } int main() { + test_get_func_attrs(); test_kernel_functor_ptr(); + test_kernel_functor_ptr_memory(); return 0; } diff --git a/sycl/test-e2e/syclcompat/kernel/Inputs/kernel_module.cpp b/sycl/test-e2e/syclcompat/kernel/Inputs/kernel_module.cpp index 85eaec5d7d632..454b06dd09064 100644 --- a/sycl/test-e2e/syclcompat/kernel/Inputs/kernel_module.cpp +++ b/sycl/test-e2e/syclcompat/kernel/Inputs/kernel_module.cpp @@ -32,20 +32,17 @@ #include -#ifdef _WIN32 -#define DLL_EXPORT __declspec(dllexport) -#else -#define DLL_EXPORT -#endif +#include void foo(int *k, sycl::nd_item<3> item_ct1, uint8_t *local_mem) { k[item_ct1.get_global_linear_id()] = item_ct1.get_global_linear_id(); } extern "C" { -DLL_EXPORT void foo_wrapper(sycl::queue &queue, const sycl::nd_range<3> &nr, - unsigned int localMemSize, void **kernelParams, - void **extra) { +SYCLCOMPAT_EXPORT void foo_wrapper(sycl::queue &queue, + const sycl::nd_range<3> &nr, + unsigned int localMemSize, + void **kernelParams, void **extra) { int *k; k = (int *)kernelParams[0]; queue.submit([&](sycl::handler &cgh) { From a5c07f53a552830c9ef28363d2b538d78f663c96 Mon Sep 17 00:00:00 2001 From: Alberto Cabrera Date: Mon, 18 Mar 2024 14:29:28 +0000 Subject: [PATCH 6/8] [SYCL][COMPAT] Fixed wrong var name for kernel params --- sycl/doc/syclcompat/README.md | 2 +- .../syclcompat/kernel/Inputs/kernel_module.cpp | 10 +++++----- 2 files changed, 6 insertions(+), 6 deletions(-) diff --git a/sycl/doc/syclcompat/README.md b/sycl/doc/syclcompat/README.md index f4eca6d756518..3ebed0d23a76c 100644 --- a/sycl/doc/syclcompat/README.md +++ b/sycl/doc/syclcompat/README.md @@ -1278,7 +1278,7 @@ static void invoke_kernel_function(kernel_function &function, sycl::range<3> group_range, sycl::range<3> local_range, unsigned int local_mem_size, - void **kernelParams, void **extra); + void **kernel_params, void **extra); } // namespace syclcompat ``` diff --git a/sycl/test-e2e/syclcompat/kernel/Inputs/kernel_module.cpp b/sycl/test-e2e/syclcompat/kernel/Inputs/kernel_module.cpp index 454b06dd09064..77afc2c15a889 100644 --- a/sycl/test-e2e/syclcompat/kernel/Inputs/kernel_module.cpp +++ b/sycl/test-e2e/syclcompat/kernel/Inputs/kernel_module.cpp @@ -41,13 +41,13 @@ void foo(int *k, sycl::nd_item<3> item_ct1, uint8_t *local_mem) { extern "C" { SYCLCOMPAT_EXPORT void foo_wrapper(sycl::queue &queue, const sycl::nd_range<3> &nr, - unsigned int localMemSize, - void **kernelParams, void **extra) { + unsigned int local_mem_size, + void **kernel_params, void **extra) { int *k; - k = (int *)kernelParams[0]; + k = (int *)kernel_params[0]; queue.submit([&](sycl::handler &cgh) { - sycl::local_accessor local_acc_ct1(sycl::range<1>(localMemSize), - cgh); + sycl::local_accessor local_acc_ct1( + sycl::range<1>(local_mem_size), cgh); cgh.parallel_for(nr, [=](sycl::nd_item<3> item_ct1) { foo(k, item_ct1, local_acc_ct1.get_multi_ptr().get()); From fbc0decb7a9f4a111cae31dfc6282eef24055875 Mon Sep 17 00:00:00 2001 From: Alberto Cabrera Date: Mon, 25 Mar 2024 10:26:16 +0000 Subject: [PATCH 7/8] [SYCL][COMPAT] Fix varname not changed in the proper header. --- sycl/include/syclcompat/kernel.hpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/sycl/include/syclcompat/kernel.hpp b/sycl/include/syclcompat/kernel.hpp index 00a9808d1d1ba..420832d561827 100644 --- a/sycl/include/syclcompat/kernel.hpp +++ b/sycl/include/syclcompat/kernel.hpp @@ -461,7 +461,7 @@ static inline void invoke_kernel_function(kernel_function &function, sycl::range<3> group_range, sycl::range<3> local_range, unsigned int local_mem_size, - void **kernelParams, void **extra) { + void **kernel_params, void **extra) { function(queue, sycl::nd_range<3>(group_range * local_range, local_range), local_mem_size, kernel_params, extra); } From 72899df864217071e19e0ab947b02e3085133651 Mon Sep 17 00:00:00 2001 From: Alberto Cabrera Date: Fri, 26 Apr 2024 12:20:22 +0100 Subject: [PATCH 8/8] [SYCL][COMPAT] Marked kernel_lin as unsupported --- sycl/test-e2e/syclcompat/kernel/kernel_lin.cpp | 2 ++ 1 file changed, 2 insertions(+) diff --git a/sycl/test-e2e/syclcompat/kernel/kernel_lin.cpp b/sycl/test-e2e/syclcompat/kernel/kernel_lin.cpp index ad341fabf9b64..eca55f738d83a 100644 --- a/sycl/test-e2e/syclcompat/kernel/kernel_lin.cpp +++ b/sycl/test-e2e/syclcompat/kernel/kernel_lin.cpp @@ -1,4 +1,6 @@ // REQUIRES: linux +// TODO: Supported for ROCM 5. Further development required to support AMDGPU. +// UNSUPPORTED: hip // RUN: %clangxx -fPIC -shared -fsycl -fsycl-targets=%{sycl_triple} %S/Inputs/kernel_module.cpp -o %t.so // RUN: %clangxx -DTEST_SHARED_LIB='"%t.so"' -ldl -fsycl -fsycl-targets=%{sycl_triple} %S/Inputs/kernel_function.cpp -o %t.out