From 58dfefb27f03f9fed3f0c087c0bf072cace128f7 Mon Sep 17 00:00:00 2001 From: Julian Oppermann Date: Mon, 17 Feb 2025 09:14:02 +0000 Subject: [PATCH 1/6] [SYCL][RTC] Query kernels by source code name Signed-off-by: Julian Oppermann --- sycl/source/detail/compiler.hpp | 2 + sycl/source/detail/device_binary_image.cpp | 1 + sycl/source/detail/device_binary_image.hpp | 4 + sycl/source/detail/jit_compiler.cpp | 17 +-- sycl/source/detail/jit_compiler.hpp | 3 +- sycl/source/detail/kernel_bundle_impl.hpp | 131 +++++++++++++----- .../kernel_compiler/kernel_compiler_sycl.cpp | 14 +- .../kernel_compiler/kernel_compiler_sycl.hpp | 4 +- .../kernel_compiler_sycl_jit.cpp | 15 +- 9 files changed, 123 insertions(+), 68 deletions(-) diff --git a/sycl/source/detail/compiler.hpp b/sycl/source/detail/compiler.hpp index 40bf97299138f..0e6a9069ce270 100644 --- a/sycl/source/detail/compiler.hpp +++ b/sycl/source/detail/compiler.hpp @@ -68,6 +68,8 @@ #define __SYCL_PROPERTY_SET_SYCL_VIRTUAL_FUNCTIONS "SYCL/virtual functions" /// PropertySetRegistry::SYCL_IMPLICIT_LOCAL_ARG defined in PropertySetIO.h #define __SYCL_PROPERTY_SET_SYCL_IMPLICIT_LOCAL_ARG "SYCL/implicit local arg" +/// PropertySetRegistry::SYCL_REGISTERED_KERNELS defined in PropertySetIO.h +#define __SYCL_PROPERTY_SET_SYCL_REGISTERED_KERNELS "SYCL/registered kernels" /// Program metadata tags recognized by the PI backends. For kernels the tag /// must appear after the kernel name. diff --git a/sycl/source/detail/device_binary_image.cpp b/sycl/source/detail/device_binary_image.cpp index 633a4269e1e78..adb55149060f2 100644 --- a/sycl/source/detail/device_binary_image.cpp +++ b/sycl/source/detail/device_binary_image.cpp @@ -195,6 +195,7 @@ void RTDeviceBinaryImage::init(sycl_device_binary Bin) { DeviceRequirements.init(Bin, __SYCL_PROPERTY_SET_SYCL_DEVICE_REQUIREMENTS); HostPipes.init(Bin, __SYCL_PROPERTY_SET_SYCL_HOST_PIPES); VirtualFunctions.init(Bin, __SYCL_PROPERTY_SET_SYCL_VIRTUAL_FUNCTIONS); + RegisteredKernels.init(Bin, __SYCL_PROPERTY_SET_SYCL_REGISTERED_KERNELS); ImageId = ImageCounter++; } diff --git a/sycl/source/detail/device_binary_image.hpp b/sycl/source/detail/device_binary_image.hpp index 50d44ad4a0e4c..211839cc41066 100644 --- a/sycl/source/detail/device_binary_image.hpp +++ b/sycl/source/detail/device_binary_image.hpp @@ -232,6 +232,9 @@ class RTDeviceBinaryImage { const PropertyRange &getHostPipes() const { return HostPipes; } const PropertyRange &getVirtualFunctions() const { return VirtualFunctions; } const PropertyRange &getImplicitLocalArg() const { return ImplicitLocalArg; } + const PropertyRange &getRegisteredKernels() const { + return RegisteredKernels; + } std::uintptr_t getImageID() const { assert(Bin && "Image ID is not available without a binary image."); @@ -258,6 +261,7 @@ class RTDeviceBinaryImage { RTDeviceBinaryImage::PropertyRange HostPipes; RTDeviceBinaryImage::PropertyRange VirtualFunctions; RTDeviceBinaryImage::PropertyRange ImplicitLocalArg; + RTDeviceBinaryImage::PropertyRange RegisteredKernels; std::vector ProgramMetadataUR; diff --git a/sycl/source/detail/jit_compiler.cpp b/sycl/source/detail/jit_compiler.cpp index 6fc88bb812a20..8046818bc3b2f 100644 --- a/sycl/source/detail/jit_compiler.cpp +++ b/sycl/source/detail/jit_compiler.cpp @@ -1249,24 +1249,11 @@ std::vector jit_compiler::encodeReqdWorkGroupSize( sycl_device_binaries jit_compiler::compileSYCL( const std::string &CompilationID, const std::string &SYCLSource, const std::vector> &IncludePairs, - const std::vector &UserArgs, std::string *LogPtr, - const std::vector &RegisteredKernelNames) { - - // RegisteredKernelNames may contain template specializations, so we just put - // them in main() which ensures they are instantiated. - std::ostringstream ss; - ss << SYCLSource << '\n'; - ss << "int main() {\n"; - for (const std::string &KernelName : RegisteredKernelNames) { - ss << " (void)" << KernelName << ";\n"; - } - ss << " return 0;\n}\n" << std::endl; - - std::string FinalSource = ss.str(); + const std::vector &UserArgs, std::string *LogPtr) { std::string SYCLFileName = CompilationID + ".cpp"; ::jit_compiler::InMemoryFile SourceFile{SYCLFileName.c_str(), - FinalSource.c_str()}; + SYCLSource.c_str()}; std::vector<::jit_compiler::InMemoryFile> IncludeFilesView; IncludeFilesView.reserve(IncludePairs.size()); diff --git a/sycl/source/detail/jit_compiler.hpp b/sycl/source/detail/jit_compiler.hpp index cf404e7bb723e..fb6b7a9b66f27 100644 --- a/sycl/source/detail/jit_compiler.hpp +++ b/sycl/source/detail/jit_compiler.hpp @@ -52,8 +52,7 @@ class jit_compiler { sycl_device_binaries compileSYCL( const std::string &CompilationID, const std::string &SYCLSource, const std::vector> &IncludePairs, - const std::vector &UserArgs, std::string *LogPtr, - const std::vector &RegisteredKernelNames); + const std::vector &UserArgs, std::string *LogPtr); void destroyDeviceBinaries(sycl_device_binaries Binaries); diff --git a/sycl/source/detail/kernel_bundle_impl.hpp b/sycl/source/detail/kernel_bundle_impl.hpp index be22afa63712a..82a3dc633a4ca 100644 --- a/sycl/source/detail/kernel_bundle_impl.hpp +++ b/sycl/source/detail/kernel_bundle_impl.hpp @@ -378,11 +378,12 @@ class kernel_bundle_impl { // oneapi_ext_kernel_compiler // program manager integration, only for sycl_jit language - kernel_bundle_impl(context Ctx, std::vector Devs, - const std::vector &KernelIDs, - std::vector KNames, - sycl_device_binaries Binaries, std::string Pfx, - syclex::source_language Lang) + kernel_bundle_impl( + context Ctx, std::vector Devs, + const std::vector &KernelIDs, + std::unordered_map &&MangledKernelNames, + sycl_device_binaries Binaries, std::string &&Prefix, + syclex::source_language Lang) : kernel_bundle_impl(std::move(Ctx), std::move(Devs), KernelIDs, bundle_state::executable) { assert(Lang == syclex::source_language::sycl_jit); @@ -392,9 +393,9 @@ class kernel_bundle_impl { // loaded via the program manager have `kernel_id`s, they can't be looked up // from the (unprefixed) kernel name. MIsInterop = true; - MKernelNames = std::move(KNames); + MMangledKernelNames = std::move(MangledKernelNames); MDeviceBinaries = Binaries; - MPrefix = std::move(Pfx); + MPrefix = std::move(Prefix); MLanguage = Lang; } @@ -501,15 +502,35 @@ class kernel_bundle_impl { // TODO: Support persistent caching. const std::string &SourceStr = std::get(MSource); + std::ostringstream SourceExt; + if (!RegisteredKernelNames.empty()) { + SourceExt << SourceStr << '\n'; + + auto EmitEntry = + [&SourceExt](const std::string &Name) -> std::ostringstream & { + SourceExt << " {\"" << Name << "\", " << Name << "}"; + return SourceExt; + }; + + SourceExt << "[[__sycl_detail__::__registered_kernels__(\n"; + for (auto It = RegisteredKernelNames.begin(), + SecondToLast = RegisteredKernelNames.end() - 1; + It != SecondToLast; ++It) { + EmitEntry(*It) << ",\n"; + } + EmitEntry(RegisteredKernelNames.back()) << "\n"; + SourceExt << ")]];\n"; + } + auto [Binaries, CompilationID] = syclex::detail::SYCL_JIT_to_SPIRV( - SourceStr, MIncludePairs, BuildOptions, LogPtr, - RegisteredKernelNames); + RegisteredKernelNames.empty() ? SourceStr : SourceExt.str(), + MIncludePairs, BuildOptions, LogPtr); auto &PM = detail::ProgramManager::getInstance(); PM.addImages(Binaries); std::vector KernelIDs; - std::vector KernelNames; + std::unordered_map MangledKernelNames; // `jit_compiler::compileSYCL(..)` uses `CompilationID + '$'` as prefix // for offload entry names. std::string Prefix = CompilationID + '$'; @@ -518,13 +539,38 @@ class kernel_bundle_impl { if (KernelName.find(Prefix) == 0) { KernelIDs.push_back(KernelID); KernelName.remove_prefix(Prefix.length()); - KernelNames.emplace_back(KernelName); + static constexpr std::string_view SYCLKernelMarker{"__sycl_kernel_"}; + if (KernelName.find(SYCLKernelMarker) == 0) { + // extern "C" declaration, register kernel without the marker. + std::string_view KernelNameWithoutMarker{KernelName}; + KernelNameWithoutMarker.remove_prefix(SYCLKernelMarker.length()); + MangledKernelNames.emplace(KernelNameWithoutMarker, KernelName); + } else { + // The marker is baked into the mangling, and we cannot easily + // adjust it. Register an identity mapping as an escape hatch. + // Users shall use `registered_kernel_names` instead, as there's + // practically no way to guess the mangled name. + MangledKernelNames.emplace(KernelName, KernelName); + } } } - return std::make_shared(MContext, MDevices, KernelIDs, - KernelNames, Binaries, Prefix, - MLanguage); + // Apply frontend information. + for (const auto *RawImg : PM.getRawDeviceImages(KernelIDs)) { + for (const sycl_device_binary_property &RKProp : + RawImg->getRegisteredKernels()) { + + auto BA = DeviceBinaryProperty(RKProp).asByteArray(); + auto MangledNameLen = BA.consume() / 8 /*bits in a byte*/; + std::string_view MangledName{ + reinterpret_cast(BA.begin()), MangledNameLen}; + MangledKernelNames.emplace(RKProp->Name, MangledName); + } + } + + return std::make_shared( + MContext, MDevices, KernelIDs, std::move(MangledKernelNames), + Binaries, std::move(Prefix), MLanguage); } ur_program_handle_t UrProgram = nullptr; @@ -642,6 +688,9 @@ class kernel_bundle_impl { } bool ext_oneapi_has_kernel(const std::string &Name) { + if (MLanguage == syclex::source_language::sycl_jit) { + return MMangledKernelNames.count(Name); + } auto it = std::find(MKernelNames.begin(), MKernelNames.end(), adjust_kernel_name(Name, MLanguage)); return it != MKernelNames.end(); @@ -650,21 +699,25 @@ class kernel_bundle_impl { kernel ext_oneapi_get_kernel(const std::string &Name, const std::shared_ptr &Self) { - if (MKernelNames.empty()) - throw sycl::exception(make_error_code(errc::invalid), - "'ext_oneapi_get_kernel' is only available in " - "kernel_bundles successfully built from " - "kernel_bundle."); + if (MLanguage == syclex::source_language::sycl_jit) { + if (MMangledKernelNames.empty()) { + throw sycl::exception( + make_error_code(errc::invalid), + "'ext_oneapi_get_kernel' is only available in kernel_bundles " + "successfully built from " + "kernel_bundle."); + } - std::string AdjustedName = adjust_kernel_name(Name, MLanguage); - if (!ext_oneapi_has_kernel(Name)) - throw sycl::exception(make_error_code(errc::invalid), - "kernel '" + AdjustedName + - "' not found in kernel_bundle"); + auto It = MMangledKernelNames.find(Name); + if (It == MMangledKernelNames.end()) { + throw sycl::exception(make_error_code(errc::invalid), + "kernel '" + Name + + "' not found in kernel_bundle"); + } - if (MLanguage == syclex::source_language::sycl_jit) { + const std::string &MangledName = It->second; auto &PM = ProgramManager::getInstance(); - auto KID = PM.getSYCLKernelID(MPrefix + AdjustedName); + auto KID = PM.getSYCLKernelID(MPrefix + MangledName); for (const auto &DevImgWithDeps : MDeviceImages) { const auto &DevImg = DevImgWithDeps.getMain(); @@ -674,7 +727,7 @@ class kernel_bundle_impl { const auto &DevImgImpl = getSyclObjImpl(DevImg); auto UrProgram = DevImgImpl->get_ur_program_ref(); auto [UrKernel, CacheMutex, ArgMask] = - PM.getOrCreateKernel(MContext, AdjustedName, + PM.getOrCreateKernel(MContext, MangledName, /*PropList=*/{}, UrProgram); auto KernelImpl = std::make_shared( UrKernel, getSyclObjImpl(MContext), DevImgImpl, Self, ArgMask, @@ -685,6 +738,18 @@ class kernel_bundle_impl { assert(false && "Malformed RTC kernel bundle"); } + if (MKernelNames.empty()) + throw sycl::exception(make_error_code(errc::invalid), + "'ext_oneapi_get_kernel' is only available in " + "kernel_bundles successfully built from " + "kernel_bundle."); + + std::string AdjustedName = adjust_kernel_name(Name, MLanguage); + if (!ext_oneapi_has_kernel(Name)) + throw sycl::exception(make_error_code(errc::invalid), + "kernel '" + AdjustedName + + "' not found in kernel_bundle"); + assert(MDeviceImages.size() > 0); const std::shared_ptr &DeviceImageImpl = detail::getSyclObjImpl(MDeviceImages[0].getMain()); @@ -877,12 +942,11 @@ class kernel_bundle_impl { } bool is_specialization_constant_set(const char *SpecName) const noexcept { - bool SetInDevImg = - std::any_of(begin(), end(), - [SpecName](const device_image_plain &DeviceImage) { - return getSyclObjImpl(DeviceImage) - ->is_specialization_constant_set(SpecName); - }); + bool SetInDevImg = std::any_of( + begin(), end(), [SpecName](const device_image_plain &DeviceImage) { + return getSyclObjImpl(DeviceImage) + ->is_specialization_constant_set(SpecName); + }); return SetInDevImg || MSpecConstValues.count(std::string{SpecName}) != 0; } @@ -973,6 +1037,7 @@ class kernel_bundle_impl { const std::variant> MSource; // only kernel_bundles created from source have KernelNames member. std::vector MKernelNames; + std::unordered_map MMangledKernelNames; sycl_device_binaries MDeviceBinaries = nullptr; std::string MPrefix; include_pairs_t MIncludePairs; diff --git a/sycl/source/detail/kernel_compiler/kernel_compiler_sycl.cpp b/sycl/source/detail/kernel_compiler/kernel_compiler_sycl.cpp index 56b79340b1309..5887cf2afcde8 100644 --- a/sycl/source/detail/kernel_compiler/kernel_compiler_sycl.cpp +++ b/sycl/source/detail/kernel_compiler/kernel_compiler_sycl.cpp @@ -303,19 +303,17 @@ bool SYCL_JIT_Compilation_Available() { #endif } -std::pair SYCL_JIT_to_SPIRV( - [[maybe_unused]] const std::string &SYCLSource, - [[maybe_unused]] include_pairs_t IncludePairs, - [[maybe_unused]] const std::vector &UserArgs, - [[maybe_unused]] std::string *LogPtr, - [[maybe_unused]] const std::vector &RegisteredKernelNames) { +std::pair +SYCL_JIT_to_SPIRV([[maybe_unused]] const std::string &SYCLSource, + [[maybe_unused]] include_pairs_t IncludePairs, + [[maybe_unused]] const std::vector &UserArgs, + [[maybe_unused]] std::string *LogPtr) { #if SYCL_EXT_JIT_ENABLE static std::atomic_uintptr_t CompilationCounter; std::string CompilationID = "rtc_" + std::to_string(CompilationCounter++); sycl_device_binaries Binaries = sycl::detail::jit_compiler::get_instance().compileSYCL( - CompilationID, SYCLSource, IncludePairs, UserArgs, LogPtr, - RegisteredKernelNames); + CompilationID, SYCLSource, IncludePairs, UserArgs, LogPtr); return std::make_pair(Binaries, std::move(CompilationID)); #else throw sycl::exception(sycl::errc::build, diff --git a/sycl/source/detail/kernel_compiler/kernel_compiler_sycl.hpp b/sycl/source/detail/kernel_compiler/kernel_compiler_sycl.hpp index 1a1a2665ae313..917259eb158e1 100644 --- a/sycl/source/detail/kernel_compiler/kernel_compiler_sycl.hpp +++ b/sycl/source/detail/kernel_compiler/kernel_compiler_sycl.hpp @@ -37,8 +37,8 @@ std::string userArgsAsString(const std::vector &UserArguments); std::pair SYCL_JIT_to_SPIRV(const std::string &Source, include_pairs_t IncludePairs, - const std::vector &UserArgs, std::string *LogPtr, - const std::vector &RegisteredKernelNames); + const std::vector &UserArgs, + std::string *LogPtr); void SYCL_JIT_destroy(sycl_device_binaries Binaries); diff --git a/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl_jit.cpp b/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl_jit.cpp index dffd1eb79c1ad..743e2689be573 100644 --- a/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl_jit.cpp +++ b/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl_jit.cpp @@ -207,14 +207,13 @@ int test_build_and_run() { // directly. sycl::kernel k = kbExe2.ext_oneapi_get_kernel("ff_cp"); - // The templated function name will have been mangled. Mapping from original - // name to mangled is not yet supported. So we cannot yet do this: - // sycl::kernel k2 = kbExe2.ext_oneapi_get_kernel("ff_templated"); - - // Instead, we can TEMPORARILY use the mangled name. Once demangling is - // supported this might no longer work. - sycl::kernel k2 = - kbExe2.ext_oneapi_get_kernel("_Z26__sycl_kernel_ff_templatedIiEvPT_S1_"); + // The templated function name will have been mangled. + sycl::kernel k2 = kbExe2.ext_oneapi_get_kernel("ff_templated"); + + // We can also use the mangled name. This escape hatch might be removed in the + // future. + assert( + kbExe2.ext_oneapi_has_kernel("_Z26__sycl_kernel_ff_templatedIiEvPT_S1_")); // Test the kernels. test_1(q, k, 37 + 5); // ff_cp seeds 37. AddEm will add 5 more. From f2de93528869ba65ecc44dca6971258b3e2e7b87 Mon Sep 17 00:00:00 2001 From: Julian Oppermann Date: Tue, 18 Feb 2025 03:25:16 +0000 Subject: [PATCH 2/6] Keep separate KernelNames list and MangledKernelNames map Signed-off-by: Julian Oppermann --- sycl/source/detail/kernel_bundle_impl.hpp | 90 ++++++++++------------- 1 file changed, 38 insertions(+), 52 deletions(-) diff --git a/sycl/source/detail/kernel_bundle_impl.hpp b/sycl/source/detail/kernel_bundle_impl.hpp index 82a3dc633a4ca..7d40bc313c03c 100644 --- a/sycl/source/detail/kernel_bundle_impl.hpp +++ b/sycl/source/detail/kernel_bundle_impl.hpp @@ -381,6 +381,7 @@ class kernel_bundle_impl { kernel_bundle_impl( context Ctx, std::vector Devs, const std::vector &KernelIDs, + std::vector &&KernelNames, std::unordered_map &&MangledKernelNames, sycl_device_binaries Binaries, std::string &&Prefix, syclex::source_language Lang) @@ -393,6 +394,7 @@ class kernel_bundle_impl { // loaded via the program manager have `kernel_id`s, they can't be looked up // from the (unprefixed) kernel name. MIsInterop = true; + MKernelNames = std::move(KernelNames); MMangledKernelNames = std::move(MangledKernelNames); MDeviceBinaries = Binaries; MPrefix = std::move(Prefix); @@ -530,6 +532,7 @@ class kernel_bundle_impl { PM.addImages(Binaries); std::vector KernelIDs; + std::vector KernelNames; std::unordered_map MangledKernelNames; // `jit_compiler::compileSYCL(..)` uses `CompilationID + '$'` as prefix // for offload entry names. @@ -539,18 +542,14 @@ class kernel_bundle_impl { if (KernelName.find(Prefix) == 0) { KernelIDs.push_back(KernelID); KernelName.remove_prefix(Prefix.length()); + KernelNames.emplace_back(KernelName); static constexpr std::string_view SYCLKernelMarker{"__sycl_kernel_"}; if (KernelName.find(SYCLKernelMarker) == 0) { - // extern "C" declaration, register kernel without the marker. + // extern "C" declaration, implicitly register kernel without the + // marker. std::string_view KernelNameWithoutMarker{KernelName}; KernelNameWithoutMarker.remove_prefix(SYCLKernelMarker.length()); MangledKernelNames.emplace(KernelNameWithoutMarker, KernelName); - } else { - // The marker is baked into the mangling, and we cannot easily - // adjust it. Register an identity mapping as an escape hatch. - // Users shall use `registered_kernel_names` instead, as there's - // practically no way to guess the mangled name. - MangledKernelNames.emplace(KernelName, KernelName); } } } @@ -569,8 +568,9 @@ class kernel_bundle_impl { } return std::make_shared( - MContext, MDevices, KernelIDs, std::move(MangledKernelNames), - Binaries, std::move(Prefix), MLanguage); + MContext, MDevices, KernelIDs, std::move(KernelNames), + std::move(MangledKernelNames), Binaries, std::move(Prefix), + MLanguage); } ur_program_handle_t UrProgram = nullptr; @@ -676,48 +676,46 @@ class kernel_bundle_impl { KernelNames, MLanguage); } - std::string adjust_kernel_name(const std::string &Name, - syclex::source_language Lang) { - // Once name demangling support is in, we won't need this. - if (Lang != syclex::source_language::sycl && - Lang != syclex::source_language::sycl_jit) - return Name; + std::string adjust_kernel_name(const std::string &Name) { + if (MLanguage == syclex::source_language::sycl_jit) { + auto It = MMangledKernelNames.find(Name); + return It == MMangledKernelNames.end() ? Name : It->second; + } + + if (MLanguage == syclex::source_language::sycl) { + bool isMangled = Name.find("__sycl_kernel_") != std::string::npos; + return isMangled ? Name : "__sycl_kernel_" + Name; + } + + return Name; + } - bool isMangled = Name.find("__sycl_kernel_") != std::string::npos; - return isMangled ? Name : "__sycl_kernel_" + Name; + bool is_kernel_name(const std::string &Name) { + return std::find(MKernelNames.begin(), MKernelNames.end(), Name) != + MKernelNames.end(); } bool ext_oneapi_has_kernel(const std::string &Name) { - if (MLanguage == syclex::source_language::sycl_jit) { - return MMangledKernelNames.count(Name); - } - auto it = std::find(MKernelNames.begin(), MKernelNames.end(), - adjust_kernel_name(Name, MLanguage)); - return it != MKernelNames.end(); + return is_kernel_name(adjust_kernel_name(Name)); } kernel ext_oneapi_get_kernel(const std::string &Name, const std::shared_ptr &Self) { - if (MLanguage == syclex::source_language::sycl_jit) { - if (MMangledKernelNames.empty()) { - throw sycl::exception( - make_error_code(errc::invalid), - "'ext_oneapi_get_kernel' is only available in kernel_bundles " - "successfully built from " - "kernel_bundle."); - } + if (MKernelNames.empty()) + throw sycl::exception(make_error_code(errc::invalid), + "'ext_oneapi_get_kernel' is only available in " + "kernel_bundles successfully built from " + "kernel_bundle."); - auto It = MMangledKernelNames.find(Name); - if (It == MMangledKernelNames.end()) { - throw sycl::exception(make_error_code(errc::invalid), - "kernel '" + Name + - "' not found in kernel_bundle"); - } + std::string AdjustedName = adjust_kernel_name(Name); + if (!is_kernel_name(AdjustedName)) + throw sycl::exception(make_error_code(errc::invalid), + "kernel '" + Name + "' not found in kernel_bundle"); - const std::string &MangledName = It->second; + if (MLanguage == syclex::source_language::sycl_jit) { auto &PM = ProgramManager::getInstance(); - auto KID = PM.getSYCLKernelID(MPrefix + MangledName); + auto KID = PM.getSYCLKernelID(MPrefix + AdjustedName); for (const auto &DevImgWithDeps : MDeviceImages) { const auto &DevImg = DevImgWithDeps.getMain(); @@ -727,7 +725,7 @@ class kernel_bundle_impl { const auto &DevImgImpl = getSyclObjImpl(DevImg); auto UrProgram = DevImgImpl->get_ur_program_ref(); auto [UrKernel, CacheMutex, ArgMask] = - PM.getOrCreateKernel(MContext, MangledName, + PM.getOrCreateKernel(MContext, AdjustedName, /*PropList=*/{}, UrProgram); auto KernelImpl = std::make_shared( UrKernel, getSyclObjImpl(MContext), DevImgImpl, Self, ArgMask, @@ -738,18 +736,6 @@ class kernel_bundle_impl { assert(false && "Malformed RTC kernel bundle"); } - if (MKernelNames.empty()) - throw sycl::exception(make_error_code(errc::invalid), - "'ext_oneapi_get_kernel' is only available in " - "kernel_bundles successfully built from " - "kernel_bundle."); - - std::string AdjustedName = adjust_kernel_name(Name, MLanguage); - if (!ext_oneapi_has_kernel(Name)) - throw sycl::exception(make_error_code(errc::invalid), - "kernel '" + AdjustedName + - "' not found in kernel_bundle"); - assert(MDeviceImages.size() > 0); const std::shared_ptr &DeviceImageImpl = detail::getSyclObjImpl(MDeviceImages[0].getMain()); From 9659324c57d0617e140ba8619f61183aa40268a0 Mon Sep 17 00:00:00 2001 From: Julian Oppermann Date: Tue, 18 Feb 2025 03:33:46 +0000 Subject: [PATCH 3/6] Clarify test Signed-off-by: Julian Oppermann --- .../test-e2e/KernelCompiler/kernel_compiler_sycl_jit.cpp | 9 ++++----- 1 file changed, 4 insertions(+), 5 deletions(-) diff --git a/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl_jit.cpp b/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl_jit.cpp index 743e2689be573..e026bbd8b7eca 100644 --- a/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl_jit.cpp +++ b/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl_jit.cpp @@ -203,15 +203,14 @@ int test_build_and_run() { syclex::properties{syclex::build_options{flags}, syclex::save_log{&log}, syclex::registered_kernel_names{"ff_templated"}}); - // extern "C" was used, so the name "ff_cp" is not mangled and can be used - // directly. + // extern "C" was used, so the name "ff_cp" is implicitly known. sycl::kernel k = kbExe2.ext_oneapi_get_kernel("ff_cp"); - // The templated function name will have been mangled. + // The templated function name was registered. sycl::kernel k2 = kbExe2.ext_oneapi_get_kernel("ff_templated"); - // We can also use the mangled name. This escape hatch might be removed in the - // future. + // We can also use the compiler-generated names directly. + assert(kbExe2.ext_oneapi_has_kernel("__sycl_kernel_ff_cp")); assert( kbExe2.ext_oneapi_has_kernel("_Z26__sycl_kernel_ff_templatedIiEvPT_S1_")); From 37c135ddb6414ed336d778cdf80384bd34ee2a28 Mon Sep 17 00:00:00 2001 From: Julian Oppermann Date: Mon, 24 Feb 2025 23:04:47 +0000 Subject: [PATCH 4/6] Revert format Signed-off-by: Julian Oppermann --- sycl/source/detail/kernel_bundle_impl.hpp | 11 ++++++----- 1 file changed, 6 insertions(+), 5 deletions(-) diff --git a/sycl/source/detail/kernel_bundle_impl.hpp b/sycl/source/detail/kernel_bundle_impl.hpp index 7d40bc313c03c..5e837e410ebb3 100644 --- a/sycl/source/detail/kernel_bundle_impl.hpp +++ b/sycl/source/detail/kernel_bundle_impl.hpp @@ -928,11 +928,12 @@ class kernel_bundle_impl { } bool is_specialization_constant_set(const char *SpecName) const noexcept { - bool SetInDevImg = std::any_of( - begin(), end(), [SpecName](const device_image_plain &DeviceImage) { - return getSyclObjImpl(DeviceImage) - ->is_specialization_constant_set(SpecName); - }); + bool SetInDevImg = + std::any_of(begin(), end(), + [SpecName](const device_image_plain &DeviceImage) { + return getSyclObjImpl(DeviceImage) + ->is_specialization_constant_set(SpecName); + }); return SetInDevImg || MSpecConstValues.count(std::string{SpecName}) != 0; } From 396ca4dc6078cfe89915984ebb868f7f094f84f5 Mon Sep 17 00:00:00 2001 From: Julian Oppermann Date: Tue, 25 Feb 2025 00:06:12 +0000 Subject: [PATCH 5/6] [SYCL][RTC] Implement `ext_oneapi_get_raw_kernel_name` Signed-off-by: Julian Oppermann --- sycl/include/sycl/kernel_bundle.hpp | 16 ++++++++++++++++ sycl/source/detail/kernel_bundle_impl.hpp | 16 ++++++++++++++++ sycl/source/kernel_bundle.cpp | 5 +++++ .../KernelCompiler/kernel_compiler_opencl.cpp | 4 ++++ .../KernelCompiler/kernel_compiler_spirv.cpp | 5 +++++ .../KernelCompiler/kernel_compiler_sycl.cpp | 5 +++++ .../KernelCompiler/kernel_compiler_sycl_jit.cpp | 11 ++++++++--- sycl/test/abi/sycl_symbols_linux.dump | 1 + 8 files changed, 60 insertions(+), 3 deletions(-) diff --git a/sycl/include/sycl/kernel_bundle.hpp b/sycl/include/sycl/kernel_bundle.hpp index 8da50f05c42a6..d77ef4319ed52 100644 --- a/sycl/include/sycl/kernel_bundle.hpp +++ b/sycl/include/sycl/kernel_bundle.hpp @@ -235,6 +235,11 @@ class __SYCL_EXPORT kernel_bundle_plain { return ext_oneapi_get_kernel(detail::string_view{name}); } + std::string ext_oneapi_get_raw_kernel_name(const std::string &name) { + return std::string{ + ext_oneapi_get_raw_kernel_name(detail::string_view{name}).c_str()}; + } + protected: // \returns a kernel object which represents the kernel identified by // kernel_id passed @@ -263,6 +268,7 @@ class __SYCL_EXPORT kernel_bundle_plain { private: bool ext_oneapi_has_kernel(detail::string_view name); kernel ext_oneapi_get_kernel(detail::string_view name); + detail::string ext_oneapi_get_raw_kernel_name(detail::string_view name); }; } // namespace detail @@ -483,6 +489,16 @@ class kernel_bundle : public detail::kernel_bundle_plain, return detail::kernel_bundle_plain::ext_oneapi_get_kernel(name); } + ///////////////////////// + // ext_oneapi_get_raw_kernel_name + // kernel_bundle must be created from source, throws if not present + ///////////////////////// + template > + std::string ext_oneapi_get_raw_kernel_name(const std::string &name) { + return detail::kernel_bundle_plain::ext_oneapi_get_raw_kernel_name(name); + } + private: kernel_bundle(detail::KernelBundleImplPtr Impl) : kernel_bundle_plain(std::move(Impl)) {} diff --git a/sycl/source/detail/kernel_bundle_impl.hpp b/sycl/source/detail/kernel_bundle_impl.hpp index b99f782b7e36f..0fdc9ff2ec874 100644 --- a/sycl/source/detail/kernel_bundle_impl.hpp +++ b/sycl/source/detail/kernel_bundle_impl.hpp @@ -748,6 +748,22 @@ class kernel_bundle_impl { return detail::createSyclObjFromImpl(KernelImpl); } + std::string ext_oneapi_get_raw_kernel_name(const std::string &Name) { + if (MKernelNames.empty()) + throw sycl::exception( + make_error_code(errc::invalid), + "'ext_oneapi_get_raw_kernel_name' is only available in " + "kernel_bundles successfully built from " + "kernel_bundle."); + + std::string AdjustedName = adjust_kernel_name(Name); + if (!is_kernel_name(AdjustedName)) + throw sycl::exception(make_error_code(errc::invalid), + "kernel '" + Name + "' not found in kernel_bundle"); + + return AdjustedName; + } + bool empty() const noexcept { return MDeviceImages.empty(); } backend get_backend() const noexcept { diff --git a/sycl/source/kernel_bundle.cpp b/sycl/source/kernel_bundle.cpp index 06a8d564221ad..19c41413ba2ed 100644 --- a/sycl/source/kernel_bundle.cpp +++ b/sycl/source/kernel_bundle.cpp @@ -137,6 +137,11 @@ kernel kernel_bundle_plain::ext_oneapi_get_kernel(detail::string_view name) { return impl->ext_oneapi_get_kernel(name.data(), impl); } +detail::string +kernel_bundle_plain::ext_oneapi_get_raw_kernel_name(detail::string_view name) { + return detail::string{impl->ext_oneapi_get_raw_kernel_name(name.data())}; +} + ////////////////////////////////// ///// sycl::detail free functions ////////////////////////////////// diff --git a/sycl/test-e2e/KernelCompiler/kernel_compiler_opencl.cpp b/sycl/test-e2e/KernelCompiler/kernel_compiler_opencl.cpp index 036bb9c86f286..8929efbc15e74 100644 --- a/sycl/test-e2e/KernelCompiler/kernel_compiler_opencl.cpp +++ b/sycl/test-e2e/KernelCompiler/kernel_compiler_opencl.cpp @@ -139,6 +139,10 @@ void test_build_and_run() { assert(hasHerKernel && "her_kernel should exist, but doesn't"); assert(!notExistKernel && "non-existing kernel should NOT exist, but does?"); + assert( + kbExe2.ext_oneapi_get_raw_kernel_name("my_kernel") == "my_kernel" && + "source code name and compiler-generated name should match, but don't"); + sycl::kernel my_kernel = kbExe2.ext_oneapi_get_kernel("my_kernel"); sycl::kernel her_kernel = kbExe2.ext_oneapi_get_kernel("her_kernel"); diff --git a/sycl/test-e2e/KernelCompiler/kernel_compiler_spirv.cpp b/sycl/test-e2e/KernelCompiler/kernel_compiler_spirv.cpp index bf6a5201708b1..e7b0e0f307b66 100644 --- a/sycl/test-e2e/KernelCompiler/kernel_compiler_spirv.cpp +++ b/sycl/test-e2e/KernelCompiler/kernel_compiler_spirv.cpp @@ -177,6 +177,11 @@ void testKernelsFromSpvFile(std::string kernels_file, sycl::queue q; auto bundle = loadKernelsFromFile(q, kernels_file); + // Test queries. + assert(bundle.ext_oneapi_has_kernel("my_kernel")); + assert(!bundle.ext_oneapi_has_kernel("not_exist")); + assert(bundle.ext_oneapi_get_raw_kernel_name("my_kernel") == "my_kernel"); + // Test simple kernel. testSimpleKernel(q, getKernel(bundle, "my_kernel"), 2, 100); diff --git a/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl.cpp b/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl.cpp index b39c3ba73d5a4..eb4e95bbddfc1 100644 --- a/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl.cpp +++ b/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl.cpp @@ -176,6 +176,11 @@ void test_build_and_run() { // extern "C" was used, so the name "ff_cp" is not mangled and can be used directly. sycl::kernel k = kbExe2.ext_oneapi_get_kernel("ff_cp"); + // Get compiler generated name and use it to query the kernel. + std::string cgn = kbExe2.ext_oneapi_get_raw_kernel_name("ff_cp"); + assert(cgn == "__sycl_kernel_ff_cp"); + assert(kbExe2.ext_oneapi_has_kernel(cgn)); + // The templated function name will have been mangled. Mapping from original // name to mangled is not yet supported. So we cannot yet do this: // sycl::kernel k2 = kbExe2.ext_oneapi_get_kernel("ff_templated"); diff --git a/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl_jit.cpp b/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl_jit.cpp index e026bbd8b7eca..f5d216f0bf89c 100644 --- a/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl_jit.cpp +++ b/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl_jit.cpp @@ -209,10 +209,15 @@ int test_build_and_run() { // The templated function name was registered. sycl::kernel k2 = kbExe2.ext_oneapi_get_kernel("ff_templated"); + // Get compiler-generated names. + std::string cgn = kbExe2.ext_oneapi_get_raw_kernel_name("ff_cp"); + std::string cgn2 = kbExe2.ext_oneapi_get_raw_kernel_name("ff_templated"); + assert(cgn == "__sycl_kernel_ff_cp"); + assert(cgn2 == "_Z26__sycl_kernel_ff_templatedIiEvPT_S1_"); + // We can also use the compiler-generated names directly. - assert(kbExe2.ext_oneapi_has_kernel("__sycl_kernel_ff_cp")); - assert( - kbExe2.ext_oneapi_has_kernel("_Z26__sycl_kernel_ff_templatedIiEvPT_S1_")); + assert(kbExe2.ext_oneapi_has_kernel(cgn)); + assert(kbExe2.ext_oneapi_has_kernel(cgn2)); // Test the kernels. test_1(q, k, 37 + 5); // ff_cp seeds 37. AddEm will add 5 more. diff --git a/sycl/test/abi/sycl_symbols_linux.dump b/sycl/test/abi/sycl_symbols_linux.dump index d762b1243217d..91df90f107ae1 100644 --- a/sycl/test/abi/sycl_symbols_linux.dump +++ b/sycl/test/abi/sycl_symbols_linux.dump @@ -3278,6 +3278,7 @@ _ZN4sycl3_V16detail18make_kernel_bundleEmRKNS0_7contextEbNS0_12bundle_stateENS0_ _ZN4sycl3_V16detail18stringifyErrorCodeEi _ZN4sycl3_V16detail19kernel_bundle_plain21ext_oneapi_get_kernelENS1_11string_viewE _ZN4sycl3_V16detail19kernel_bundle_plain21ext_oneapi_has_kernelENS1_11string_viewE +_ZN4sycl3_V16detail19kernel_bundle_plain30ext_oneapi_get_raw_kernel_nameENS1_11string_viewE _ZN4sycl3_V16detail19kernel_bundle_plain32set_specialization_constant_implEPKcPvm _ZN4sycl3_V16detail20associateWithHandlerERNS0_7handlerEPNS1_16AccessorBaseHostENS0_6access6targetE _ZN4sycl3_V16detail20associateWithHandlerERNS0_7handlerEPNS1_28SampledImageAccessorBaseHostENS0_12image_targetE From 30879de882d25a0fa35de921aceb1703c246ad76 Mon Sep 17 00:00:00 2001 From: Julian Oppermann Date: Wed, 26 Feb 2025 09:00:16 +0000 Subject: [PATCH 6/6] Add Windows symbols Signed-off-by: Julian Oppermann --- sycl/test/abi/sycl_symbols_windows.dump | 2 ++ 1 file changed, 2 insertions(+) diff --git a/sycl/test/abi/sycl_symbols_windows.dump b/sycl/test/abi/sycl_symbols_windows.dump index 9704acd6c643f..46888a69df725 100644 --- a/sycl/test/abi/sycl_symbols_windows.dump +++ b/sycl/test/abi/sycl_symbols_windows.dump @@ -3878,6 +3878,8 @@ ?ext_oneapi_get_kernel@kernel_bundle_plain@detail@_V1@sycl@@QEAA?AVkernel@34@AEBV?$basic_string@DU?$char_traits@D@std@@V?$allocator@D@2@@std@@@Z ?ext_oneapi_get_last_event@queue@_V1@sycl@@QEBA?AV?$optional@Vevent@_V1@sycl@@@std@@XZ ?ext_oneapi_get_last_event_impl@queue@_V1@sycl@@AEBA?AV?$optional@Vevent@_V1@sycl@@@detail@23@XZ +?ext_oneapi_get_raw_kernel_name@kernel_bundle_plain@detail@_V1@sycl@@AEAA?AVstring@234@Vstring_view@234@@Z +?ext_oneapi_get_raw_kernel_name@kernel_bundle_plain@detail@_V1@sycl@@QEAA?AV?$basic_string@DU?$char_traits@D@std@@V?$allocator@D@2@@std@@AEBV56@@Z ?ext_oneapi_get_state@queue@_V1@sycl@@QEBA?AW4queue_state@experimental@oneapi@ext@23@XZ ?ext_oneapi_graph@handler@_V1@sycl@@QEAAXV?$command_graph@$00@experimental@oneapi@ext@23@@Z ?ext_oneapi_graph@queue@_V1@sycl@@QEAA?AVevent@23@V?$command_graph@$00@experimental@oneapi@ext@23@AEBUcode_location@detail@23@@Z