From c309d1ccd2cf4bca13829e223097499ebfb4c701 Mon Sep 17 00:00:00 2001 From: KornevNikita Date: Wed, 23 Nov 2022 07:06:48 -0800 Subject: [PATCH 1/5] [SYCL] Implement sycl::is_compatible() function --- sycl/include/sycl/kernel_bundle.hpp | 3 ++- sycl/source/kernel_bundle.cpp | 32 +++++++++++++++++++++++++++++ 2 files changed, 34 insertions(+), 1 deletion(-) diff --git a/sycl/include/sycl/kernel_bundle.hpp b/sycl/include/sycl/kernel_bundle.hpp index 23b540cc810b7..791eee545052c 100644 --- a/sycl/include/sycl/kernel_bundle.hpp +++ b/sycl/include/sycl/kernel_bundle.hpp @@ -564,7 +564,8 @@ bool has_kernel_bundle(const context &Ctx, const std::vector &Devs) { /// \returns true if all of the kernels identified by KernelIDs are compatible /// with the device Dev. -bool is_compatible(const std::vector &KernelIDs, const device &Dev); +__SYCL_EXPORT bool is_compatible(const std::vector &KernelIDs, + const device &Dev); template bool is_compatible(const device &Dev) { return is_compatible({get_kernel_id()}, Dev); diff --git a/sycl/source/kernel_bundle.cpp b/sycl/source/kernel_bundle.cpp index db94b79561a7e..1ec3f502ded06 100644 --- a/sycl/source/kernel_bundle.cpp +++ b/sycl/source/kernel_bundle.cpp @@ -8,6 +8,7 @@ #include #include +#include #include @@ -294,5 +295,36 @@ std::vector get_kernel_ids() { return detail::ProgramManager::getInstance().getAllSYCLKernelIDs(); } +bool is_compatible(const std::vector &KernelIDs, const device &Dev) { + for (const auto &KernelId : KernelIDs) { + const detail::RTDeviceBinaryImage &Img = + detail::ProgramManager::getInstance().getDeviceImage( + detail::OSUtil::ExeModuleHandle, KernelId.get_name(), context(Dev), + Dev); + const detail::RTDeviceBinaryImage::PropertyRange &ARange = + Img.getDeviceRequirements(); + for (detail::RTDeviceBinaryImage::PropertyRange::ConstIterator It : + ARange) { + using namespace std::literals; + if ((*It)->Name != "aspects"sv) + continue; + detail::ByteArray Aspects = + detail::DeviceBinaryProperty(*It).asByteArray(); + // 8 because we need to skip 64-bits of size of the byte array + auto *AIt = reinterpret_cast(&Aspects[8]); + auto *AEnd = + reinterpret_cast(&Aspects[0] + Aspects.size()); + while (AIt != AEnd) { + auto Aspect = static_cast(*AIt); + if (!Dev.has(Aspect)) + return false; + ++AIt; + } + } + } + + return true; +} + } // __SYCL_INLINE_VER_NAMESPACE(_V1) } // namespace sycl From 22800fa01b9825e1633d14d669f6dd5b3d2c7681 Mon Sep 17 00:00:00 2001 From: KornevNikita Date: Wed, 23 Nov 2022 08:43:15 -0800 Subject: [PATCH 2/5] Test fix --- sycl/source/kernel_bundle.cpp | 2 ++ sycl/test/abi/sycl_symbols_linux.dump | 1 + 2 files changed, 3 insertions(+) diff --git a/sycl/source/kernel_bundle.cpp b/sycl/source/kernel_bundle.cpp index 1ec3f502ded06..b4f4ac61cdefc 100644 --- a/sycl/source/kernel_bundle.cpp +++ b/sycl/source/kernel_bundle.cpp @@ -296,6 +296,8 @@ std::vector get_kernel_ids() { } bool is_compatible(const std::vector &KernelIDs, const device &Dev) { + if (KernelIDs.empty()) + return false; for (const auto &KernelId : KernelIDs) { const detail::RTDeviceBinaryImage &Img = detail::ProgramManager::getInstance().getDeviceImage( diff --git a/sycl/test/abi/sycl_symbols_linux.dump b/sycl/test/abi/sycl_symbols_linux.dump index 82348e967c551..4b38a06742ea1 100644 --- a/sycl/test/abi/sycl_symbols_linux.dump +++ b/sycl/test/abi/sycl_symbols_linux.dump @@ -3619,6 +3619,7 @@ _ZN4sycl3_V113aligned_allocEmmRKNS0_5queueENS0_3usm5allocERKNS0_13property_listE _ZN4sycl3_V113aligned_allocEmmRKNS0_5queueENS0_3usm5allocERKNS0_6detail13code_locationE _ZN4sycl3_V113aligned_allocEmmRKNS0_6deviceERKNS0_7contextENS0_3usm5allocERKNS0_13property_listERKNS0_6detail13code_locationE _ZN4sycl3_V113aligned_allocEmmRKNS0_6deviceERKNS0_7contextENS0_3usm5allocERKNS0_6detail13code_locationE +_ZN4sycl3_V113is_compatibleERKSt6vectorINS0_9kernel_idESaIS2_EERKNS0_6deviceE _ZN4sycl3_V113malloc_deviceEmRKNS0_5queueERKNS0_13property_listERKNS0_6detail13code_locationE _ZN4sycl3_V113malloc_deviceEmRKNS0_5queueERKNS0_6detail13code_locationE _ZN4sycl3_V113malloc_deviceEmRKNS0_6deviceERKNS0_7contextERKNS0_13property_listERKNS0_6detail13code_locationE From 1e9ea5e091ae2df93aefebd837b616b54fdfbe54 Mon Sep 17 00:00:00 2001 From: KornevNikita Date: Thu, 24 Nov 2022 06:26:17 -0800 Subject: [PATCH 3/5] Apply suggestions --- sycl/source/kernel_bundle.cpp | 15 ++++++--------- 1 file changed, 6 insertions(+), 9 deletions(-) diff --git a/sycl/source/kernel_bundle.cpp b/sycl/source/kernel_bundle.cpp index b4f4ac61cdefc..ee29610b37349 100644 --- a/sycl/source/kernel_bundle.cpp +++ b/sycl/source/kernel_bundle.cpp @@ -12,6 +12,8 @@ #include +#include + namespace sycl { __SYCL_INLINE_VER_NAMESPACE(_V1) { @@ -296,8 +298,6 @@ std::vector get_kernel_ids() { } bool is_compatible(const std::vector &KernelIDs, const device &Dev) { - if (KernelIDs.empty()) - return false; for (const auto &KernelId : KernelIDs) { const detail::RTDeviceBinaryImage &Img = detail::ProgramManager::getInstance().getDeviceImage( @@ -312,15 +312,12 @@ bool is_compatible(const std::vector &KernelIDs, const device &Dev) { continue; detail::ByteArray Aspects = detail::DeviceBinaryProperty(*It).asByteArray(); - // 8 because we need to skip 64-bits of size of the byte array - auto *AIt = reinterpret_cast(&Aspects[8]); - auto *AEnd = - reinterpret_cast(&Aspects[0] + Aspects.size()); - while (AIt != AEnd) { - auto Aspect = static_cast(*AIt); + // Drop 8 bytes describing the size of the byte array + Aspects.dropBytes(8); + while (!Aspects.empty()) { + aspect Aspect = Aspects.consume(); if (!Dev.has(Aspect)) return false; - ++AIt; } } } From f3339307ab771bb38b15c76094c809bfff2d910b Mon Sep 17 00:00:00 2001 From: KornevNikita Date: Thu, 24 Nov 2022 06:27:12 -0800 Subject: [PATCH 4/5] Typo --- sycl/source/kernel_bundle.cpp | 2 -- 1 file changed, 2 deletions(-) diff --git a/sycl/source/kernel_bundle.cpp b/sycl/source/kernel_bundle.cpp index ee29610b37349..826170316f124 100644 --- a/sycl/source/kernel_bundle.cpp +++ b/sycl/source/kernel_bundle.cpp @@ -12,8 +12,6 @@ #include -#include - namespace sycl { __SYCL_INLINE_VER_NAMESPACE(_V1) { From 8464034b944a709c3750f72b417122703f3dcbe4 Mon Sep 17 00:00:00 2001 From: "Kornev, Nikita" Date: Thu, 24 Nov 2022 18:15:02 +0100 Subject: [PATCH 5/5] Adjust sycl_symbols_windows --- sycl/test/abi/sycl_symbols_windows.dump | 1 + 1 file changed, 1 insertion(+) diff --git a/sycl/test/abi/sycl_symbols_windows.dump b/sycl/test/abi/sycl_symbols_windows.dump index b1b892c6e24ac..36657b27b85ce 100644 --- a/sycl/test/abi/sycl_symbols_windows.dump +++ b/sycl/test/abi/sycl_symbols_windows.dump @@ -948,6 +948,7 @@ ?isValidModeForSourceAccessor@handler@_V1@sycl@@CA_NW4mode@access@23@@Z ?isValidTargetForExplicitOp@handler@_V1@sycl@@CA_NW4target@access@23@@Z ?is_accelerator@device@_V1@sycl@@QEBA_NXZ +?is_compatible@_V1@sycl@@YA_NAEBV?$vector@Vkernel_id@_V1@sycl@@V?$allocator@Vkernel_id@_V1@sycl@@@std@@@std@@AEBVdevice@12@@Z ?is_cpu@device@_V1@sycl@@QEBA_NXZ ?is_gpu@device@_V1@sycl@@QEBA_NXZ ?is_host@context@_V1@sycl@@QEBA_NXZ