diff --git a/sycl/include/CL/sycl/detail/service_kernel_names.hpp b/sycl/include/CL/sycl/detail/service_kernel_names.hpp new file mode 100644 index 0000000000000..86a5defeae8c4 --- /dev/null +++ b/sycl/include/CL/sycl/detail/service_kernel_names.hpp @@ -0,0 +1,21 @@ +//==-------- service_kernels.hpp - SYCL service kernel name types ----------==// +// +// 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 +// +//===----------------------------------------------------------------------===// + +#pragma once + +__SYCL_INLINE_NAMESPACE(cl) { +namespace sycl { +namespace detail { +namespace __sycl_service_kernel__ { + +class AssertInfoCopier; + +} // namespace __sycl_service_kernel__ +} // namespace detail +} // namespace sycl +} // __SYCL_INLINE_NAMESPACE(cl) diff --git a/sycl/include/CL/sycl/queue.hpp b/sycl/include/CL/sycl/queue.hpp index 01f6cbc99015e..11db9e697a82c 100644 --- a/sycl/include/CL/sycl/queue.hpp +++ b/sycl/include/CL/sycl/queue.hpp @@ -12,6 +12,7 @@ #include #include #include +#include #include #include #include @@ -79,11 +80,10 @@ class queue; namespace detail { class queue_impl; #if __SYCL_USE_FALLBACK_ASSERT -class AssertInfoCopier; static event submitAssertCapture(queue &, event &, queue *, const detail::code_location &); #endif -} +} // namespace detail /// Encapsulates a single SYCL queue which schedules kernels on a SYCL device. /// @@ -1167,7 +1167,7 @@ event submitAssertCapture(queue &Self, event &Event, queue *SecondaryQueue, auto Acc = Buffer.get_access(CGH); - CGH.single_task([Acc] { + CGH.single_task<__sycl_service_kernel__::AssertInfoCopier>([Acc] { #ifdef __SYCL_DEVICE_ONLY__ __devicelib_assert_read(&Acc[0]); #else diff --git a/sycl/source/detail/program_manager/program_manager.cpp b/sycl/source/detail/program_manager/program_manager.cpp index b916039bda9ee..b9b05adb8eb2e 100644 --- a/sycl/source/detail/program_manager/program_manager.cpp +++ b/sycl/source/detail/program_manager/program_manager.cpp @@ -1044,6 +1044,16 @@ void ProgramManager::addImages(pi_device_binaries DeviceBinary) { auto Result = KSIdMap.insert(std::make_pair(EntriesIt->name, KSId)); (void)Result; assert(Result.second && "Kernel sets are not disjoint"); + + // Skip creating unique kernel ID if it is a service kernel. + // SYCL service kernels are identified by having + // __sycl_service_kernel__ in the mangled name, primarily as part of + // the namespace of the name type. + if (std::strstr(EntriesIt->name, "__sycl_service_kernel__")) { + m_ServiceKernels.insert(EntriesIt->name); + continue; + } + // ... and create a unique kernel ID for the entry std::shared_ptr KernelIDImpl = std::make_shared(EntriesIt->name); @@ -1323,7 +1333,6 @@ ProgramManager::getSYCLDeviceImagesWithCompatibleState( if (!compatibleWithDevice(BinImage, Dev)) continue; - // TODO: Cache kernel_ids std::vector KernelIDs; // Collect kernel names for the image pi_device_binary DevBin = @@ -1333,11 +1342,23 @@ ProgramManager::getSYCLDeviceImagesWithCompatibleState( for (_pi_offload_entry EntriesIt = DevBin->EntriesBegin; EntriesIt != DevBin->EntriesEnd; ++EntriesIt) { auto KernelID = m_KernelIDs.find(EntriesIt->name); - assert(KernelID != m_KernelIDs.end() && - "Kernel ID in device binary missing from cache"); + + if (KernelID == m_KernelIDs.end()) { + // Service kernels do not have kernel IDs + assert(m_ServiceKernels.find(EntriesIt->name) != + m_ServiceKernels.end() && + "Kernel ID in device binary missing from cache"); + continue; + } + KernelIDs.push_back(KernelID->second); } } + + // If the image does not contain any non-service kernels we can skip it. + if (KernelIDs.empty()) + continue; + // device_image_impl expects kernel ids to be sorted for fast search std::sort(KernelIDs.begin(), KernelIDs.end(), LessByNameComp{}); diff --git a/sycl/source/detail/program_manager/program_manager.hpp b/sycl/source/detail/program_manager/program_manager.hpp index 04c7202499721..8fd1278d386b0 100644 --- a/sycl/source/detail/program_manager/program_manager.hpp +++ b/sycl/source/detail/program_manager/program_manager.hpp @@ -22,6 +22,7 @@ #include #include #include +#include #include // +++ Entry points referenced by the offload wrapper object { @@ -275,7 +276,7 @@ class ProgramManager { /// Maps names of kernels to their unique kernel IDs. /// TODO: Use std::unordered_set with transparent hash and equality functions /// when C++20 is enabled for the runtime library. - /// Access must be guarded by the m_KernelIDsMutex mutex + /// Access must be guarded by the m_KernelIDsMutex mutex. std::unordered_map m_KernelIDs; /// Protects kernel ID cache. @@ -284,6 +285,14 @@ class ProgramManager { /// \ref Sync::getGlobalLock() while holding this mutex. std::mutex m_KernelIDsMutex; + /// Caches all found service kernels to expedite future checks. A SYCL service + /// kernel is a kernel that has not been defined by the user but is instead + /// generated by the SYCL runtime. Service kernel name types must be declared + /// in the sycl::detail::__sycl_service_kernel__ namespace which is + /// exclusively used for this purpose. + /// Access must be guarded by the m_KernelIDsMutex mutex. + std::unordered_set m_ServiceKernels; + // Keeps track of pi_program to image correspondence. Needed for: // - knowing which specialization constants are used in the program and // injecting their current values before compiling the SPIR-V; the binary diff --git a/sycl/unittests/SYCL2020/KernelID.cpp b/sycl/unittests/SYCL2020/KernelID.cpp index fc4d34f99ac3e..baf3ef4bc7796 100644 --- a/sycl/unittests/SYCL2020/KernelID.cpp +++ b/sycl/unittests/SYCL2020/KernelID.cpp @@ -17,6 +17,7 @@ class TestKernel1; class TestKernel2; class TestKernel3; +class ServiceKernel1; __SYCL_INLINE_NAMESPACE(cl) { namespace sycl { @@ -57,6 +58,19 @@ template <> struct KernelInfo { static constexpr bool callsAnyThisFreeFunction() { return false; } }; +template <> struct KernelInfo { + static constexpr unsigned getNumParams() { return 0; } + static const kernel_param_desc_t &getParamDesc(int) { + static kernel_param_desc_t Dummy; + return Dummy; + } + static constexpr const char *getName() { + return "_ZTSN2cl4sycl6detail23__sycl_service_kernel__14ServiceKernel1"; + } + static constexpr bool isESIMD() { return false; } + static constexpr bool callsThisItem() { return false; } + static constexpr bool callsAnyThisFreeFunction() { return false; } +}; } // namespace detail } // namespace sycl } // __SYCL_INLINE_NAMESPACE(cl) @@ -84,7 +98,9 @@ generateDefaultImage(std::initializer_list Kernels) { static sycl::unittest::PiImage Imgs[2] = { generateDefaultImage({"KernelID_TestKernel1", "KernelID_TestKernel3"}), - generateDefaultImage({"KernelID_TestKernel2"})}; + generateDefaultImage( + {"KernelID_TestKernel2", + "_ZTSN2cl4sycl6detail23__sycl_service_kernel__14ServiceKernel1"})}; static sycl::unittest::PiImageArray<2> ImgArray{Imgs}; TEST(KernelID, AllProgramKernelIds) { @@ -106,6 +122,20 @@ TEST(KernelID, AllProgramKernelIds) { } } +TEST(KernelID, NoServiceKernelIds) { + const char *ServiceKernel1Name = + sycl::detail::KernelInfo::getName(); + + std::vector AllKernelIDs = sycl::get_kernel_ids(); + + auto NoFoundServiceKernelID = std::none_of( + AllKernelIDs.begin(), AllKernelIDs.end(), [=](sycl::kernel_id KernelID) { + return strcmp(KernelID.get_name(), ServiceKernel1Name) == 0; + }); + + EXPECT_TRUE(NoFoundServiceKernelID); +} + TEST(KernelID, FreeKernelIDEqualsKernelBundleId) { sycl::platform Plt{sycl::default_selector()}; if (Plt.is_host()) { diff --git a/sycl/unittests/assert/assert.cpp b/sycl/unittests/assert/assert.cpp index 4d583d7588610..b2f0e3acdc934 100644 --- a/sycl/unittests/assert/assert.cpp +++ b/sycl/unittests/assert/assert.cpp @@ -51,9 +51,10 @@ template <> struct KernelInfo { static constexpr const kernel_param_desc_t Signatures[] = { {kernel_param_kind_t::kind_accessor, 4062, 0}}; -template <> struct KernelInfo<::sycl::detail::AssertInfoCopier> { +template <> +struct KernelInfo<::sycl::detail::__sycl_service_kernel__::AssertInfoCopier> { static constexpr const char *getName() { - return "_ZTSN2cl4sycl6detail16AssertInfoCopierE"; + return "_ZTSN2cl4sycl6detail23__sycl_service_kernel__16AssertInfoCopierE"; } static constexpr unsigned getNumParams() { return 1; } static constexpr const kernel_param_desc_t &getParamDesc(unsigned Idx) { @@ -73,7 +74,7 @@ static sycl::unittest::PiImage generateDefaultImage() { static const std::string KernelName = "TestKernel"; static const std::string CopierKernelName = - "_ZTSN2cl4sycl6detail16AssertInfoCopierE"; + "_ZTSN2cl4sycl6detail23__sycl_service_kernel__16AssertInfoCopierE"; PiPropertySet PropSet; @@ -98,7 +99,7 @@ static sycl::unittest::PiImage generateCopierKernelImage() { using namespace sycl::unittest; static const std::string CopierKernelName = - "_ZTSN2cl4sycl6detail16AssertInfoCopierE"; + "_ZTSN2cl4sycl6detail23__sycl_service_kernel__16AssertInfoCopierE"; PiPropertySet PropSet; @@ -391,3 +392,17 @@ TEST(Assert, TestPositive) { } #endif // _WIN32 } + +TEST(Assert, TestAssertServiceKernelHidden) { + const char *AssertServiceKernelName = sycl::detail::KernelInfo< + sycl::detail::__sycl_service_kernel__::AssertInfoCopier>::getName(); + + std::vector AllKernelIDs = sycl::get_kernel_ids(); + + auto NoFoundServiceKernelID = std::none_of( + AllKernelIDs.begin(), AllKernelIDs.end(), [=](sycl::kernel_id KernelID) { + return strcmp(KernelID.get_name(), AssertServiceKernelName) == 0; + }); + + EXPECT_TRUE(NoFoundServiceKernelID); +}