From c947aec0aa27d956ab0cc10fe4927d9e9b1cc05f Mon Sep 17 00:00:00 2001 From: Sergey Semenov Date: Thu, 24 Jul 2025 10:08:34 -0700 Subject: [PATCH 1/2] [SYCL] Fix adjusted kernel name handling in preview builds Remove usage of a temporary string where the program manager receives and stores a string view in the preview mode. --- sycl/source/detail/device_image_impl.cpp | 8 ++++---- sycl/source/detail/device_image_impl.hpp | 18 ++++++++++++------ sycl/source/detail/kernel_bundle_impl.hpp | 2 +- .../detail/program_manager/program_manager.cpp | 8 +++++--- 4 files changed, 22 insertions(+), 14 deletions(-) diff --git a/sycl/source/detail/device_image_impl.cpp b/sycl/source/detail/device_image_impl.cpp index 20b7465f30851..fa7f3f41f8f03 100644 --- a/sycl/source/detail/device_image_impl.cpp +++ b/sycl/source/detail/device_image_impl.cpp @@ -21,18 +21,18 @@ std::shared_ptr device_image_impl::tryGetExtensionKernel( !((getOriginMask() & ImageOriginSYCLBIN) && hasKernelName(Name))) return nullptr; - std::string AdjustedName = adjustKernelName(Name); + std::string_view AdjustedName = adjustKernelName(Name); if (MRTCBinInfo && MRTCBinInfo->MLanguage == syclex::source_language::sycl) { auto &PM = ProgramManager::getInstance(); for (const std::string &Prefix : MRTCBinInfo->MPrefixes) { - auto KID = PM.tryGetSYCLKernelID(Prefix + AdjustedName); + auto KID = PM.tryGetSYCLKernelID(Prefix + std::string(AdjustedName)); if (!KID || !has_kernel(*KID)) continue; auto UrProgram = get_ur_program(); auto [UrKernel, CacheMutex, ArgMask] = - PM.getOrCreateKernel(Context, AdjustedName, + PM.getOrCreateKernel(Context, KernelNameStrT(AdjustedName), /*PropList=*/{}, UrProgram); return std::make_shared(UrKernel, *getSyclObjImpl(Context), shared_from_this(), OwnerBundle, @@ -44,7 +44,7 @@ std::shared_ptr device_image_impl::tryGetExtensionKernel( ur_program_handle_t UrProgram = get_ur_program(); detail::adapter_impl &Adapter = getSyclObjImpl(Context)->getAdapter(); ur_kernel_handle_t UrKernel = nullptr; - Adapter.call(UrProgram, AdjustedName.c_str(), + Adapter.call(UrProgram, AdjustedName.data(), &UrKernel); // Kernel created by urKernelCreate is implicitly retained. diff --git a/sycl/source/detail/device_image_impl.hpp b/sycl/source/detail/device_image_impl.hpp index 97daad803229d..379b10c37d222 100644 --- a/sycl/source/detail/device_image_impl.hpp +++ b/sycl/source/detail/device_image_impl.hpp @@ -152,9 +152,12 @@ class ManagedDeviceBinaries { sycl_device_binaries MBinaries; }; +// Using ordered containers for heterogenous lookup. +// TODO change to unordered containers after switching to C++20. using MangledKernelNameMapT = std::map>; using KernelNameSetT = std::set>; -using KernelNameToArgMaskMap = std::unordered_map; +using KernelNameToArgMaskMap = + std::map>; // Information unique to images compiled at runtime through the // ext_oneapi_kernel_compiler extension. @@ -619,17 +622,20 @@ class device_image_impl #pragma warning(pop) #endif - std::string adjustKernelName(std::string_view Name) const { + std::string_view adjustKernelName(std::string_view Name) const { if (MOrigins & ImageOriginSYCLBIN) { constexpr std::string_view KernelPrefix = "__sycl_kernel_"; if (Name.size() > KernelPrefix.size() && Name.substr(0, KernelPrefix.size()) == KernelPrefix) - return Name.data(); - return std::string{KernelPrefix} + Name.data(); + return Name; + auto It = + MKernelNames.find(std::string(KernelPrefix) + std::string(Name)); + assert(It != MKernelNames.end() && "Adjusted name not found"); + return *It; } if (!MRTCBinInfo.has_value()) - return Name.data(); + return Name; if (MRTCBinInfo->MLanguage == syclex::source_language::sycl) { auto It = MRTCBinInfo->MMangledKernelNames.find(Name); @@ -637,7 +643,7 @@ class device_image_impl return It->second; } - return Name.data(); + return Name; } bool hasKernelName(std::string_view Name) const { diff --git a/sycl/source/detail/kernel_bundle_impl.hpp b/sycl/source/detail/kernel_bundle_impl.hpp index 0cd49a7afcadc..f27bbee00c586 100644 --- a/sycl/source/detail/kernel_bundle_impl.hpp +++ b/sycl/source/detail/kernel_bundle_impl.hpp @@ -692,7 +692,7 @@ class kernel_bundle_impl throw sycl::exception(make_error_code(errc::invalid), "kernel '" + Name + "' not found in kernel_bundle"); - return It->adjustKernelName(Name); + return std::string(It->adjustKernelName(Name)); } bool ext_oneapi_has_device_global(const std::string &Name) const { diff --git a/sycl/source/detail/program_manager/program_manager.cpp b/sycl/source/detail/program_manager/program_manager.cpp index b2eb455726a7e..b599f49adae69 100644 --- a/sycl/source/detail/program_manager/program_manager.cpp +++ b/sycl/source/detail/program_manager/program_manager.cpp @@ -2825,7 +2825,7 @@ ProgramManager::compile(const DevImgPlainWithDeps &ImgWithDeps, setSpecializationConstants(InputImpl, Prog, Adapter); KernelNameSetT KernelNames = InputImpl.getKernelNames(); - std::unordered_map EliminatedKernelArgMasks = + std::map> EliminatedKernelArgMasks = InputImpl.getEliminatedKernelArgMasks(); std::optional RTCInfo = @@ -3015,7 +3015,8 @@ ProgramManager::link(const std::vector &Imgs, RTCInfoPtrs; RTCInfoPtrs.reserve(Imgs.size()); KernelNameSetT MergedKernelNames; - std::unordered_map MergedEliminatedKernelArgMasks; + std::map> + MergedEliminatedKernelArgMasks; for (const device_image_plain &DevImg : Imgs) { device_image_impl &DevImgImpl = *getSyclObjImpl(DevImg); CombinedOrigins |= DevImgImpl.getOriginMask(); @@ -3099,7 +3100,8 @@ ProgramManager::build(const DevImgPlainWithDeps &DevImgWithDeps, RTCInfoPtrs; RTCInfoPtrs.reserve(DevImgWithDeps.size()); KernelNameSetT MergedKernelNames; - std::unordered_map MergedEliminatedKernelArgMasks; + std::map> + MergedEliminatedKernelArgMasks; for (const device_image_plain &DevImg : DevImgWithDeps) { device_image_impl &DevImgImpl = *getSyclObjImpl(DevImg); RTCInfoPtrs.emplace_back(&(DevImgImpl.getRTCInfo())); From 397ef3ee4ef712b61a321d58af7ec3ad883d96ab Mon Sep 17 00:00:00 2001 From: Sergey Semenov Date: Fri, 25 Jul 2025 09:19:54 -0700 Subject: [PATCH 2/2] Fix hasKernelName with SYCLBIN --- sycl/source/detail/device_image_impl.cpp | 2 +- sycl/source/detail/device_image_impl.hpp | 61 +++++++++++++++-------- sycl/source/detail/kernel_bundle_impl.hpp | 2 +- 3 files changed, 41 insertions(+), 24 deletions(-) diff --git a/sycl/source/detail/device_image_impl.cpp b/sycl/source/detail/device_image_impl.cpp index fa7f3f41f8f03..80222f9bd0ec0 100644 --- a/sycl/source/detail/device_image_impl.cpp +++ b/sycl/source/detail/device_image_impl.cpp @@ -21,7 +21,7 @@ std::shared_ptr device_image_impl::tryGetExtensionKernel( !((getOriginMask() & ImageOriginSYCLBIN) && hasKernelName(Name))) return nullptr; - std::string_view AdjustedName = adjustKernelName(Name); + std::string_view AdjustedName = getAdjustedKernelNameStrView(Name); if (MRTCBinInfo && MRTCBinInfo->MLanguage == syclex::source_language::sycl) { auto &PM = ProgramManager::getInstance(); for (const std::string &Prefix : MRTCBinInfo->MPrefixes) { diff --git a/sycl/source/detail/device_image_impl.hpp b/sycl/source/detail/device_image_impl.hpp index 379b10c37d222..f11782237db9b 100644 --- a/sycl/source/detail/device_image_impl.hpp +++ b/sycl/source/detail/device_image_impl.hpp @@ -622,35 +622,21 @@ class device_image_impl #pragma warning(pop) #endif - std::string_view adjustKernelName(std::string_view Name) const { - if (MOrigins & ImageOriginSYCLBIN) { - constexpr std::string_view KernelPrefix = "__sycl_kernel_"; - if (Name.size() > KernelPrefix.size() && - Name.substr(0, KernelPrefix.size()) == KernelPrefix) - return Name; - auto It = - MKernelNames.find(std::string(KernelPrefix) + std::string(Name)); - assert(It != MKernelNames.end() && "Adjusted name not found"); - return *It; - } - - if (!MRTCBinInfo.has_value()) - return Name; - - if (MRTCBinInfo->MLanguage == syclex::source_language::sycl) { - auto It = MRTCBinInfo->MMangledKernelNames.find(Name); - if (It != MRTCBinInfo->MMangledKernelNames.end()) - return It->second; - } + // Assumes the kernel is contained within this image. + std::string_view getAdjustedKernelNameStrView(std::string_view Name) const { + return getAdjustedKernelNameImpl(Name); + } - return Name; + std::string getAdjustedKernelNameStr(std::string_view Name) const { + return getAdjustedKernelNameImpl(Name); } bool hasKernelName(std::string_view Name) const { return (getOriginMask() & (ImageOriginKernelCompiler | ImageOriginSYCLBIN)) && !Name.empty() && - MKernelNames.find(adjustKernelName(Name)) != MKernelNames.end(); + MKernelNames.find(getAdjustedKernelNameStr(Name)) != + MKernelNames.end(); } std::shared_ptr @@ -846,6 +832,37 @@ class device_image_impl } private: + template + RetT getAdjustedKernelNameImpl(std::string_view Name) const { + if (MOrigins & ImageOriginSYCLBIN) { + constexpr std::string_view KernelPrefix = "__sycl_kernel_"; + if (Name.size() > KernelPrefix.size() && + Name.substr(0, KernelPrefix.size()) == KernelPrefix) + return RetT(Name); + std::string AdjustedNameStr = + std::string(KernelPrefix) + std::string(Name); + if constexpr (std::is_same_v) { + return AdjustedNameStr; + } else { + static_assert(std::is_same_v); + auto It = MKernelNames.find(AdjustedNameStr); + assert(It != MKernelNames.end() && "Adjusted name not found"); + return *It; + } + } + + if (!MRTCBinInfo.has_value()) + return RetT(Name); + + if (MRTCBinInfo->MLanguage == syclex::source_language::sycl) { + auto It = MRTCBinInfo->MMangledKernelNames.find(Name); + if (It != MRTCBinInfo->MMangledKernelNames.end()) + return It->second; + } + + return RetT(Name); + } + bool hasRTDeviceBinaryImage() const noexcept { return std::holds_alternative(MBinImage) && get_bin_image_ref() != nullptr; diff --git a/sycl/source/detail/kernel_bundle_impl.hpp b/sycl/source/detail/kernel_bundle_impl.hpp index f27bbee00c586..5e54fc54f96af 100644 --- a/sycl/source/detail/kernel_bundle_impl.hpp +++ b/sycl/source/detail/kernel_bundle_impl.hpp @@ -692,7 +692,7 @@ class kernel_bundle_impl throw sycl::exception(make_error_code(errc::invalid), "kernel '" + Name + "' not found in kernel_bundle"); - return std::string(It->adjustKernelName(Name)); + return It->getAdjustedKernelNameStr(Name); } bool ext_oneapi_has_device_global(const std::string &Name) const {