diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 7a0b2dede0a7..27bdf6966e9c 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -1111,18 +1111,24 @@ static std::pair constructFreeFunctionKernelName( SemaSYCL &SemaSYCLRef, const FunctionDecl *FreeFunc, MangleContext &MC) { SmallString<256> Result; llvm::raw_svector_ostream Out(Result); + std::string NewName; std::string StableName; - MC.mangleName(FreeFunc, Out); - std::string MangledName(Out.str()); - size_t StartNums = MangledName.find_first_of("0123456789"); - size_t EndNums = MangledName.find_first_not_of("0123456789", StartNums); - size_t NameLength = - std::stoi(MangledName.substr(StartNums, EndNums - StartNums)); - size_t NewNameLength = 14 /*length of __sycl_kernel_*/ + NameLength; - std::string NewName = MangledName.substr(0, StartNums) + - std::to_string(NewNameLength) + "__sycl_kernel_" + - MangledName.substr(EndNums); + // Handle extern "C" + if (FreeFunc->getLanguageLinkage() == CLanguageLinkage) { + const IdentifierInfo *II = FreeFunc->getIdentifier(); + NewName = "__sycl_kernel_" + II->getName().str(); + } else { + MC.mangleName(FreeFunc, Out); + std::string MangledName(Out.str()); + size_t StartNums = MangledName.find_first_of("0123456789"); + size_t EndNums = MangledName.find_first_not_of("0123456789", StartNums); + size_t NameLength = + std::stoi(MangledName.substr(StartNums, EndNums - StartNums)); + size_t NewNameLength = 14 /*length of __sycl_kernel_*/ + NameLength; + NewName = MangledName.substr(0, StartNums) + std::to_string(NewNameLength) + + "__sycl_kernel_" + MangledName.substr(EndNums); + } StableName = NewName; return {NewName, StableName}; } diff --git a/sycl/include/sycl/ext/oneapi/properties/property.hpp b/sycl/include/sycl/ext/oneapi/properties/property.hpp index 3f1bb28268d3..50b385f3f0ca 100644 --- a/sycl/include/sycl/ext/oneapi/properties/property.hpp +++ b/sycl/include/sycl/ext/oneapi/properties/property.hpp @@ -209,8 +209,10 @@ enum PropKind : uint32_t { CallsIndirectly = 68, InputDataPlacement = 69, OutputDataPlacement = 70, + IncludeFiles = 71, + RegisteredKernelNames = 72, // PropKindSize must always be the last value. - PropKindSize = 71, + PropKindSize = 73, }; struct property_key_base_tag {}; diff --git a/sycl/include/sycl/kernel_bundle.hpp b/sycl/include/sycl/kernel_bundle.hpp index 5bba4735561a..99663183ad8d 100644 --- a/sycl/include/sycl/kernel_bundle.hpp +++ b/sycl/include/sycl/kernel_bundle.hpp @@ -814,6 +814,32 @@ build(const kernel_bundle &InputBundle, namespace ext::oneapi::experimental { +namespace detail { +struct create_bundle_from_source_props; +struct build_source_bundle_props; +} // namespace detail + +///////////////////////// +// PropertyT syclex::include_files +///////////////////////// +struct include_files + : detail::run_time_property_key { + include_files(); + include_files(const std::string &name, const std::string &content) { + record.emplace_back(std::make_pair(name, content)); + } + void add(const std::string &name, const std::string &content) { + record.emplace_back(std::make_pair(name, content)); + } + std::vector> record; +}; +using include_files_key = include_files; + +template <> +struct is_property_key_of + : std::true_type {}; + ///////////////////////// // PropertyT syclex::build_options ///////////////////////// @@ -826,8 +852,7 @@ struct build_options using build_options_key = build_options; template <> -struct is_property_key_of> +struct is_property_key_of : std::true_type {}; ///////////////////////// @@ -840,72 +865,132 @@ struct save_log : detail::run_time_property_key { using save_log_key = save_log; template <> -struct is_property_key_of> +struct is_property_key_of : std::true_type {}; ///////////////////////// -// syclex::is_source_kernel_bundle_supported +// PropertyT syclex::registered_kernel_names ///////////////////////// +struct registered_kernel_names + : detail::run_time_property_key { + std::vector kernel_names; + registered_kernel_names() {} + registered_kernel_names(const std::string &knArg) : kernel_names{knArg} {} + registered_kernel_names(const std::vector &knsArg) + : kernel_names(knsArg) {} + void add(const std::string &name) { kernel_names.push_back(name); } +}; +using registered_kernel_names_key = registered_kernel_names; + +template <> +struct is_property_key_of : std::true_type { +}; + +namespace detail { +// forward decls __SYCL_EXPORT bool is_source_kernel_bundle_supported(backend BE, source_language Language); -///////////////////////// -// syclex::create_kernel_bundle_from_source -///////////////////////// - __SYCL_EXPORT kernel_bundle -create_kernel_bundle_from_source(const context &SyclContext, - source_language Language, - const std::string &Source); +make_kernel_bundle_from_source( + const context &SyclContext, source_language Language, + const std::string &Source, + std::vector> IncludePairsVec); #if (!defined(_HAS_STD_BYTE) || _HAS_STD_BYTE != 0) __SYCL_EXPORT kernel_bundle -create_kernel_bundle_from_source(const context &SyclContext, - source_language Language, - const std::vector &Bytes); +make_kernel_bundle_from_source( + const context &SyclContext, source_language Language, + const std::vector &Bytes, + std::vector> IncludePairsVec); #endif -///////////////////////// -// syclex::build(source_kb) => exe_kb -///////////////////////// -namespace detail { -// forward decl __SYCL_EXPORT kernel_bundle build_from_source(kernel_bundle &SourceKB, const std::vector &Devices, const std::vector &BuildOptions, - std::string *LogPtr); + std::string *LogPtr, + const std::vector &RegisteredKernelNames); } // namespace detail +///////////////////////// +// syclex::create_kernel_bundle_from_source +///////////////////////// +template < + typename PropertyListT = empty_properties_t, + typename = std::enable_if_t< + is_property_list_v && + detail::all_props_are_keys_of::value>> +kernel_bundle create_kernel_bundle_from_source( + const context &SyclContext, source_language Language, + const std::string &Source, PropertyListT props = {}) { + std::vector> IncludePairsVec; + if constexpr (props.template has_property()) { + IncludePairsVec = props.template get_property().record; + } + + return detail::make_kernel_bundle_from_source(SyclContext, Language, Source, + IncludePairsVec); +} + +#if (!defined(_HAS_STD_BYTE) || _HAS_STD_BYTE != 0) +template < + typename PropertyListT = empty_properties_t, + typename = std::enable_if_t< + is_property_list_v && + detail::all_props_are_keys_of::value>> +kernel_bundle create_kernel_bundle_from_source( + const context &SyclContext, source_language Language, + const std::vector &Bytes, PropertyListT props = {}) { + std::vector> IncludePairsVec; + if constexpr (props.template has_property()) { + IncludePairsVec = props.template get_property().record; + } + + return detail::make_kernel_bundle_from_source(SyclContext, Language, Bytes, + IncludePairsVec); +} +#endif + +///////////////////////// +// syclex::build(source_kb) => exe_kb +///////////////////////// + template && - detail::all_props_are_keys_of< - kernel_bundle, - PropertyListT>::value>> + detail::all_props_are_keys_of::value>> kernel_bundle build(kernel_bundle &SourceKB, const std::vector &Devices, PropertyListT props = {}) { std::vector BuildOptionsVec; std::string *LogPtr = nullptr; + std::vector RegisteredKernelNamesVec; if constexpr (props.template has_property()) { BuildOptionsVec = props.template get_property().opts; } if constexpr (props.template has_property()) { LogPtr = props.template get_property().log; } - return detail::build_from_source(SourceKB, Devices, BuildOptionsVec, LogPtr); + if constexpr (props.template has_property()) { + RegisteredKernelNamesVec = + props.template get_property().kernel_names; + } + return detail::build_from_source(SourceKB, Devices, BuildOptionsVec, LogPtr, + RegisteredKernelNamesVec); } template && - detail::all_props_are_keys_of< - kernel_bundle, - PropertyListT>::value>> + detail::all_props_are_keys_of::value>> kernel_bundle build(kernel_bundle &SourceKB, PropertyListT props = {}) { diff --git a/sycl/include/sycl/kernel_bundle_enums.hpp b/sycl/include/sycl/kernel_bundle_enums.hpp index 936b0de3879f..fd53f8cd3a74 100644 --- a/sycl/include/sycl/kernel_bundle_enums.hpp +++ b/sycl/include/sycl/kernel_bundle_enums.hpp @@ -20,7 +20,7 @@ enum class bundle_state : char { namespace ext::oneapi::experimental { -enum class source_language : int { opencl = 0, spirv = 1 /* sycl, cuda */ }; +enum class source_language : int { opencl = 0, spirv = 1, sycl = 2 /* cuda */ }; // opencl versions struct cl_version { diff --git a/sycl/source/CMakeLists.txt b/sycl/source/CMakeLists.txt index f915ef4e2cb8..e72cdebf506c 100644 --- a/sycl/source/CMakeLists.txt +++ b/sycl/source/CMakeLists.txt @@ -217,6 +217,7 @@ set(SYCL_COMMON_SOURCES "detail/jit_compiler.cpp" "detail/jit_device_binaries.cpp" "detail/kernel_compiler/kernel_compiler_opencl.cpp" + "detail/kernel_compiler/kernel_compiler_sycl.cpp" "detail/kernel_impl.cpp" "detail/kernel_program_cache.cpp" "detail/memory_manager.cpp" diff --git a/sycl/source/detail/device_impl.cpp b/sycl/source/detail/device_impl.cpp index 8547a40d4b99..3f0d44b90a79 100644 --- a/sycl/source/detail/device_impl.cpp +++ b/sycl/source/detail/device_impl.cpp @@ -862,7 +862,8 @@ bool device_impl::isGetDeviceAndHostTimerSupported() { bool device_impl::extOneapiCanCompile( ext::oneapi::experimental::source_language Language) { try { - return is_source_kernel_bundle_supported(getBackend(), Language); + return sycl::ext::oneapi::experimental::detail:: + is_source_kernel_bundle_supported(getBackend(), Language); } catch (sycl::exception &) { return false; } diff --git a/sycl/source/detail/kernel_bundle_impl.hpp b/sycl/source/detail/kernel_bundle_impl.hpp index 55586b6d2b5a..35a5b690ccb8 100644 --- a/sycl/source/detail/kernel_bundle_impl.hpp +++ b/sycl/source/detail/kernel_bundle_impl.hpp @@ -10,6 +10,7 @@ #include #include +#include #include #include #include @@ -329,12 +330,15 @@ class kernel_bundle_impl { } } + using include_pairs_t = + std::vector>; // oneapi_ext_kernel_compiler // construct from source string kernel_bundle_impl(const context &Context, syclex::source_language Lang, - const std::string &Src) + const std::string &Src, include_pairs_t IncludePairsVec) : MContext(Context), MDevices(Context.get_devices()), - MState(bundle_state::ext_oneapi_source), Language(Lang), Source(Src) {} + MState(bundle_state::ext_oneapi_source), Language(Lang), Source(Src), + IncludePairs(IncludePairsVec) {} // oneapi_ext_kernel_compiler // construct from source bytes @@ -348,16 +352,19 @@ class kernel_bundle_impl { // interop constructor kernel_bundle_impl(context Ctx, std::vector Devs, device_image_plain &DevImage, - std::vector KNames) + std::vector KNames, + syclex::source_language Lang) : kernel_bundle_impl(Ctx, Devs, DevImage) { MState = bundle_state::executable; KernelNames = KNames; + Language = Lang; } std::shared_ptr build_from_source(const std::vector Devices, const std::vector &BuildOptions, - std::string *LogPtr) { + std::string *LogPtr, + const std::vector &RegisteredKernelNames) { assert(MState == bundle_state::ext_oneapi_source && "bundle_state::ext_oneapi_source required"); @@ -397,6 +404,12 @@ class kernel_bundle_impl { [](std::byte B) { return static_cast(B); }); return Result; } + if (Language == syclex::source_language::sycl) { + const auto &SourceStr = std::get(this->Source); + return syclex::detail::SYCL_to_SPIRV(SourceStr, IncludePairs, + BuildOptions, LogPtr, + RegisteredKernelNames); + } throw sycl::exception( make_error_code(errc::invalid), "OpenCL C and SPIR-V are the only supported languages at this time"); @@ -437,11 +450,22 @@ class kernel_bundle_impl { PiProgram); device_image_plain DevImg{DevImgImpl}; return std::make_shared(MContext, MDevices, DevImg, - KernelNames); + KernelNames, Language); + } + + 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) + return Name; + + bool isMangled = Name.find("__sycl_kernel_") != std::string::npos; + return isMangled ? Name : "__sycl_kernel_" + Name; } bool ext_oneapi_has_kernel(const std::string &Name) { - auto it = std::find(KernelNames.begin(), KernelNames.end(), Name); + auto it = std::find(KernelNames.begin(), KernelNames.end(), + adjust_kernel_name(Name, Language)); return it != KernelNames.end(); } @@ -454,9 +478,11 @@ class kernel_bundle_impl { "kernel_bundles successfully built from " "kernel_bundle."); + std::string AdjustedName = adjust_kernel_name(Name, Language); if (!ext_oneapi_has_kernel(Name)) throw sycl::exception(make_error_code(errc::invalid), - "kernel '" + Name + "' not found in kernel_bundle"); + "kernel '" + AdjustedName + + "' not found in kernel_bundle"); assert(MDeviceImages.size() > 0); const std::shared_ptr &DeviceImageImpl = @@ -465,7 +491,8 @@ class kernel_bundle_impl { ContextImplPtr ContextImpl = getSyclObjImpl(MContext); const PluginPtr &Plugin = ContextImpl->getPlugin(); sycl::detail::pi::PiKernel PiKernel = nullptr; - Plugin->call(PiProgram, Name.c_str(), &PiKernel); + Plugin->call(PiProgram, AdjustedName.c_str(), + &PiKernel); // Kernel created by piKernelCreate is implicitly retained. std::shared_ptr KernelImpl = std::make_shared( @@ -710,11 +737,14 @@ class kernel_bundle_impl { SpecConstMapT MSpecConstValues; bool MIsInterop = false; bundle_state MState; - // ext_oneapi_kernel_compiler : Source, Languauge, KernelNames - const syclex::source_language Language = syclex::source_language::opencl; + + // ext_oneapi_kernel_compiler : Source, Languauge, KernelNames, IncludePairs + // Language is for both state::source and state::executable. + syclex::source_language Language = syclex::source_language::opencl; const std::variant> Source; // only kernel_bundles created from source have KernelNames member. std::vector KernelNames; + include_pairs_t IncludePairs; }; } // namespace detail diff --git a/sycl/source/detail/kernel_compiler/kernel_compiler_sycl.cpp b/sycl/source/detail/kernel_compiler/kernel_compiler_sycl.cpp new file mode 100644 index 000000000000..98f296bb6ceb --- /dev/null +++ b/sycl/source/detail/kernel_compiler/kernel_compiler_sycl.cpp @@ -0,0 +1,265 @@ +//==-- kernel_compiler_opencl.cpp OpenCL kernel compilation support -==// +// +// 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 +// +//===----------------------------------------------------------------------===// + +#include "kernel_compiler_sycl.hpp" +#include // make_error_code + +#if __GNUC__ && __GNUC__ < 8 + +// std::filesystem is not availalbe for GCC < 8 +// and much of the cross-platform file handling code depends upon it. +// Given that this extension is experimental and that the file +// handling aspects are most likely temporary, it makes sense to +// simply not support GCC<8. + +namespace sycl { +inline namespace _V1 { +namespace ext::oneapi::experimental { +namespace detail { + +bool SYCL_Compilation_Available() { return false; } + +spirv_vec_t +SYCL_to_SPIRV(const std::string &SYCLSource, include_pairs_t IncludePairs, + const std::vector &UserArgs, std::string *LogPtr, + const std::vector &RegisteredKernelNames) { + throw sycl::exception(sycl::errc::build, + "kernel_compiler does not support GCC<8"); +} +} // namespace detail +} // namespace ext::oneapi::experimental +} // namespace _V1 +} // namespace sycl + +#else + +#include +#include +#include +#include +#include +#include + +namespace sycl { +inline namespace _V1 { +namespace ext::oneapi::experimental { +namespace detail { + +std::string generateSemiUniqueId() { + // Get the current time as a time_t object. + std::time_t CurrentTime = std::time(nullptr); + + // Convert time_t to a string with format YYYYMMDD_HHMMSS. + std::tm *LocalTime = std::localtime(&CurrentTime); + std::stringstream Ss; + Ss << std::put_time(LocalTime, "%Y%m%d_%H%M%S"); + + // Amend with random number. + std::random_device Rd; + int RandomNumber = Rd() % 900 + 100; + Ss << "_" << std::setfill('0') << std::setw(3) << RandomNumber; + + return Ss.str(); +} + +std::filesystem::path prepareWS(const std::string &Id) { + const std::filesystem::path TmpDirectoryPath = + std::filesystem::temp_directory_path(); + std::filesystem::path NewDirectoryPath = TmpDirectoryPath / Id; + + try { + std::filesystem::create_directories(NewDirectoryPath); + } catch (const std::filesystem::filesystem_error &E) { + throw sycl::exception(sycl::errc::build, E.what()); + } + + return NewDirectoryPath; +} + +std::string userArgsAsString(const std::vector &UserArguments) { + return std::accumulate(UserArguments.begin(), UserArguments.end(), + std::string(""), + [](const std::string &A, const std::string &B) { + return A.empty() ? B : A + " " + B; + }); +} + +void outputPreamble(std::ofstream &Os, const std::filesystem::path &FilePath, + const std::string &Id, + const std::vector &UserArgs) { + + Os << "/*\n"; + Os << " clang++ -fsycl -o " << Id << ".bin "; + Os << userArgsAsString(UserArgs); + Os << " -fno-sycl-dead-args-optimization -fsycl-dump-device-code=./ " << Id; + Os << ".cpp \n */" << std::endl; +} + +std::filesystem::path +outputCpp(const std::filesystem::path &ParentDir, const std::string &Id, + std::string RawCodeString, const std::vector &UserArgs, + const std::vector &RegisteredKernelNames) { + std::filesystem::path FilePath = ParentDir / (Id + ".cpp"); + std::ofstream Outfile(FilePath, std::ios::out | std::ios::trunc); + + if (Outfile.is_open()) { + outputPreamble(Outfile, FilePath, Id, UserArgs); + Outfile << RawCodeString << std::endl; + + // Temporarily needed until -c works with -fsycl-dump-spirv. + Outfile << "int main() {\n"; + for (const std::string &KernelName : RegisteredKernelNames) { + Outfile << " " << KernelName << ";\n"; + } + Outfile << " return 0;\n}\n" << std::endl; + + Outfile.close(); + } else { + throw sycl::exception(sycl::errc::build, + "Failed to open .cpp file for write: " + + FilePath.string()); + } + + return FilePath; +} + +void outputIncludeFiles(const std::filesystem::path &Dirpath, + include_pairs_t IncludePairs) { + using pairStrings = std::pair; + for (pairStrings p : IncludePairs) { + std::filesystem::path FilePath = Dirpath / p.first; + std::ofstream outfile(FilePath, std::ios::out | std::ios::trunc); + if (outfile.is_open()) { + outfile << p.second << std::endl; + + outfile.close(); + } else { + throw sycl::exception(sycl::errc::build, + "Failed to open include file for write: " + + FilePath.string()); + } + } +} + +std::string getCompilerName() { +#ifdef __WIN32 + std::string Compiler = "clang++.exe"; +#else + std::string Compiler = "clang++"; +#endif + return Compiler; +} + +void invokeCompiler(const std::filesystem::path &FPath, + const std::filesystem::path &DPath, const std::string &Id, + const std::vector &UserArgs, + std::string *LogPtr) { + + std::filesystem::path FilePath(FPath); + std::filesystem::path ParentDir(DPath); + std::filesystem::path TargetPath = ParentDir / (Id + ".bin"); + std::filesystem::path LogPath = ParentDir / "compilation_log.txt"; + std::string Compiler = getCompilerName(); + + std::string Command = + Compiler + " -fsycl -o " + TargetPath.make_preferred().string() + " " + + userArgsAsString(UserArgs) + + " -fno-sycl-dead-args-optimization -fsycl-dump-device-code=" + + ParentDir.make_preferred().string() + " " + + FilePath.make_preferred().string() + " 2> " + + LogPath.make_preferred().string(); + + int Result = std::system(Command.c_str()); + + // Read the log file contents into the log variable. + std::string CompileLog; + std::ifstream LogStream; + LogStream.open(LogPath); + if (LogStream.is_open()) { + std::stringstream LogBuffer; + LogBuffer << LogStream.rdbuf(); + CompileLog.append(LogBuffer.str()); + if (LogPtr != nullptr) + LogPtr->append(LogBuffer.str()); + + } else if (Result == 0 && LogPtr != nullptr) { + // If there was a compilation problem, we want to report that (below), + // not a mere "missing log" error. + throw sycl::exception(sycl::errc::build, + "failure retrieving compilation log"); + } + + if (Result != 0) { + throw sycl::exception(sycl::errc::build, + "Compile failure: " + std::to_string(Result) + " " + + CompileLog); + } +} + +std::filesystem::path findSpv(const std::filesystem::path &ParentDir, + const std::string &Id) { + std::regex PatternRegex(Id + R"(.*\.spv)"); + + // Iterate through all files in the directory matching the pattern. + for (const auto &Entry : std::filesystem::directory_iterator(ParentDir)) { + if (Entry.is_regular_file() && + std::regex_match(Entry.path().filename().string(), PatternRegex)) { + return Entry.path(); // Return the path if it matches the SPV pattern. + } + } + + // File not found, throw. + throw sycl::exception(sycl::errc::build, "SPIRV output matching " + Id + + " missing from " + + ParentDir.filename().string()); +} + +spirv_vec_t loadSpvFromFile(const std::filesystem::path &FileName) { + std::ifstream SpvStream(FileName, std::ios::binary); + SpvStream.seekg(0, std::ios::end); + size_t Size = SpvStream.tellg(); + SpvStream.seekg(0); + spirv_vec_t Spv(Size); + SpvStream.read(reinterpret_cast(Spv.data()), Size); + + return Spv; +} + +spirv_vec_t +SYCL_to_SPIRV(const std::string &SYCLSource, include_pairs_t IncludePairs, + const std::vector &UserArgs, std::string *LogPtr, + const std::vector &RegisteredKernelNames) { + // clang-format off + const std::string id = generateSemiUniqueId(); + const std::filesystem::path ParentDir = prepareWS(id); + std::filesystem::path FilePath = outputCpp(ParentDir, id, SYCLSource, UserArgs, RegisteredKernelNames); + outputIncludeFiles(ParentDir, IncludePairs); + invokeCompiler(FilePath, ParentDir, id, UserArgs, LogPtr); + std::filesystem::path SpvPath = findSpv(ParentDir, id); + return loadSpvFromFile(SpvPath); + // clang-format on +} + +bool SYCL_Compilation_Available() { + // Is compiler on $PATH ? We try to invoke it. + std::string id = generateSemiUniqueId(); + const std::filesystem::path tmp = std::filesystem::temp_directory_path(); + std::filesystem::path DumpPath = tmp / (id + "_version.txt"); + std::string Compiler = getCompilerName(); + std::string TestCommand = + Compiler + " --version &> " + DumpPath.make_preferred().string(); + int result = std::system(TestCommand.c_str()); + + return (result == 0); +} + +} // namespace detail +} // namespace ext::oneapi::experimental +} // namespace _V1 +} // namespace sycl +#endif diff --git a/sycl/source/detail/kernel_compiler/kernel_compiler_sycl.hpp b/sycl/source/detail/kernel_compiler/kernel_compiler_sycl.hpp new file mode 100644 index 000000000000..dfff9ac839e8 --- /dev/null +++ b/sycl/source/detail/kernel_compiler/kernel_compiler_sycl.hpp @@ -0,0 +1,38 @@ +//==-- kernel_compiler_sycl.hpp SYCL kernel compilation support -==// +// +// 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 + +#include +#include // __SYCL_EXPORT +#include + +#include // std::accumulate +#include +#include + +namespace sycl { +inline namespace _V1 { +namespace ext::oneapi::experimental { +namespace detail { + +using spirv_vec_t = std::vector; +using include_pairs_t = std::vector>; + +spirv_vec_t +SYCL_to_SPIRV(const std::string &Source, include_pairs_t IncludePairs, + const std::vector &UserArgs, std::string *LogPtr, + const std::vector &RegisteredKernelNames); + +bool SYCL_Compilation_Available(); + +} // namespace detail +} // namespace ext::oneapi::experimental + +} // namespace _V1 +} // namespace sycl diff --git a/sycl/source/kernel_bundle.cpp b/sycl/source/kernel_bundle.cpp index aace54af59ac..ee02cb6d1b93 100644 --- a/sycl/source/kernel_bundle.cpp +++ b/sycl/source/kernel_bundle.cpp @@ -9,6 +9,7 @@ #include #include #include +#include #include #include @@ -362,19 +363,22 @@ using source_kb = kernel_bundle; using exe_kb = kernel_bundle; using kernel_bundle_impl = sycl::detail::kernel_bundle_impl; +namespace detail { + ///////////////////////// -// syclex::is_source_kernel_bundle_supported +// syclex::detail::is_source_kernel_bundle_supported ///////////////////////// bool is_source_kernel_bundle_supported(backend BE, source_language Language) { // Support is limited to the opencl and level_zero backends. bool BE_Acceptable = (BE == sycl::backend::ext_oneapi_level_zero) || (BE == sycl::backend::opencl); if (BE_Acceptable) { - // At the moment, OpenCL and SPIR-V are the only supported languages. if (Language == source_language::opencl) { return detail::OpenCLC_Compilation_Available(); } else if (Language == source_language::spirv) { return true; + } else if (Language == source_language::sycl) { + return detail::SYCL_Compilation_Available(); } } @@ -383,12 +387,15 @@ bool is_source_kernel_bundle_supported(backend BE, source_language Language) { } ///////////////////////// -// syclex::create_kernel_bundle_from_source +// syclex::detail::create_kernel_bundle_from_source ///////////////////////// -source_kb create_kernel_bundle_from_source(const context &SyclContext, - source_language Language, - const std::string &Source) { +using include_pairs_t = std::vector>; + +source_kb make_kernel_bundle_from_source(const context &SyclContext, + source_language Language, + const std::string &Source, + include_pairs_t IncludePairs) { // TODO: if we later support a "reason" why support isn't present // (like a missing shared library etc.) it'd be nice to include it in // the exception message here. @@ -397,20 +404,27 @@ source_kb create_kernel_bundle_from_source(const context &SyclContext, throw sycl::exception(make_error_code(errc::invalid), "kernel_bundle creation from source not supported"); + // throw if include not supported? awaiting guidance + // if(!IncludePairs.empty() && is_include_supported(Languuage)){ throw invalid + // } + std::shared_ptr KBImpl = - std::make_shared(SyclContext, Language, Source); + std::make_shared(SyclContext, Language, Source, + IncludePairs); return sycl::detail::createSyclObjFromImpl(KBImpl); } -source_kb -create_kernel_bundle_from_source(const context &SyclContext, - source_language Language, - const std::vector &Bytes) { +source_kb make_kernel_bundle_from_source(const context &SyclContext, + source_language Language, + const std::vector &Bytes, + include_pairs_t IncludePairs) { backend BE = SyclContext.get_backend(); if (!is_source_kernel_bundle_supported(BE, Language)) throw sycl::exception(make_error_code(errc::invalid), "kernel_bundle creation from source not supported"); + // throw if !IncludePairs.empty() ? awaiting guidance. + std::shared_ptr KBImpl = std::make_shared(SyclContext, Language, Bytes); return sycl::detail::createSyclObjFromImpl(KBImpl); @@ -419,17 +433,17 @@ create_kernel_bundle_from_source(const context &SyclContext, ///////////////////////// // syclex::detail::build_from_source(source_kb) => exe_kb ///////////////////////// -namespace detail { -exe_kb build_from_source(source_kb &SourceKB, - const std::vector &Devices, - const std::vector &BuildOptions, - std::string *LogPtr) { +exe_kb +build_from_source(source_kb &SourceKB, const std::vector &Devices, + const std::vector &BuildOptions, + std::string *LogPtr, + const std::vector &RegisteredKernelNames) { std::vector UniqueDevices = sycl::detail::removeDuplicateDevices(Devices); std::shared_ptr sourceImpl = getSyclObjImpl(SourceKB); - std::shared_ptr KBImpl = - sourceImpl->build_from_source(UniqueDevices, BuildOptions, LogPtr); + std::shared_ptr KBImpl = sourceImpl->build_from_source( + UniqueDevices, BuildOptions, LogPtr, RegisteredKernelNames); return sycl::detail::createSyclObjFromImpl(KBImpl); } diff --git a/sycl/test-e2e/EnqueueFunctions/kernel_shortcut_with_kb.cpp b/sycl/test-e2e/EnqueueFunctions/kernel_shortcut_with_kb.cpp index 7176def8ec7d..485266c8a3d0 100644 --- a/sycl/test-e2e/EnqueueFunctions/kernel_shortcut_with_kb.cpp +++ b/sycl/test-e2e/EnqueueFunctions/kernel_shortcut_with_kb.cpp @@ -21,8 +21,8 @@ constexpr size_t N = 1024; int main() { sycl::queue Q; - if (!oneapiext::is_source_kernel_bundle_supported( - Q.get_backend(), oneapiext::source_language::opencl)) { + if (!Q.get_device().ext_oneapi_can_compile( + oneapiext::source_language::opencl)) { std::cout << "Backend does not support OpenCL C source kernel bundle extension: " << Q.get_backend() << std::endl; diff --git a/sycl/test-e2e/EnqueueFunctions/kernel_submit_with_event_and_kb.cpp b/sycl/test-e2e/EnqueueFunctions/kernel_submit_with_event_and_kb.cpp index 0e8988574fd5..b2731740a60e 100644 --- a/sycl/test-e2e/EnqueueFunctions/kernel_submit_with_event_and_kb.cpp +++ b/sycl/test-e2e/EnqueueFunctions/kernel_submit_with_event_and_kb.cpp @@ -21,8 +21,8 @@ int main() { sycl::queue Q; int Memory[N] = {0}; - if (!oneapiext::is_source_kernel_bundle_supported( - Q.get_backend(), oneapiext::source_language::opencl)) { + if (!Q.get_device().ext_oneapi_can_compile( + oneapiext::source_language::opencl)) { std::cout << "Backend does not support OpenCL C source kernel bundle extension: " << Q.get_backend() << std::endl; diff --git a/sycl/test-e2e/EnqueueFunctions/kernel_submit_with_kb.cpp b/sycl/test-e2e/EnqueueFunctions/kernel_submit_with_kb.cpp index be29a73b87ee..2651a803b509 100644 --- a/sycl/test-e2e/EnqueueFunctions/kernel_submit_with_kb.cpp +++ b/sycl/test-e2e/EnqueueFunctions/kernel_submit_with_kb.cpp @@ -21,8 +21,8 @@ int main() { sycl::queue Q; int Memory[N] = {0}; - if (!oneapiext::is_source_kernel_bundle_supported( - Q.get_backend(), oneapiext::source_language::opencl)) { + if (!Q.get_device().ext_oneapi_can_compile( + oneapiext::source_language::opencl)) { std::cout << "Backend does not support OpenCL C source kernel bundle extension: " << Q.get_backend() << std::endl; diff --git a/sycl/test-e2e/KernelCompiler/kernel_compiler.cpp b/sycl/test-e2e/KernelCompiler/kernel_compiler_opencl.cpp similarity index 92% rename from sycl/test-e2e/KernelCompiler/kernel_compiler.cpp rename to sycl/test-e2e/KernelCompiler/kernel_compiler_opencl.cpp index 356a26f7be2e..79b72ee19b0a 100644 --- a/sycl/test-e2e/KernelCompiler/kernel_compiler.cpp +++ b/sycl/test-e2e/KernelCompiler/kernel_compiler_opencl.cpp @@ -1,4 +1,4 @@ -//==- kernel_compiler.cpp --- kernel_compiler extension tests -------------==// +//==- kernel_compiler_opencl.cpp --- kernel_compiler extension tests -------==// // // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. // See https://llvm.org/LICENSE.txt for license information. @@ -85,12 +85,13 @@ void test_build_and_run() { sycl::context ctx{d}; sycl::queue q{ctx, d}; - bool ok = syclex::is_source_kernel_bundle_supported( - ctx.get_backend(), syclex::source_language::opencl); + bool ok = + q.get_device().ext_oneapi_can_compile(syclex::source_language::opencl); if (!ok) { - std::cout << "Apparently this backend does not support OpenCL C source " + std::cout << "Apparently this device does not support OpenCL C source " "kernel bundle extension: " - << ctx.get_backend() << std::endl; + << q.get_device().get_info() + << std::endl; return; } @@ -141,8 +142,8 @@ void test_error() { sycl::context ctx{d}; sycl::queue q{ctx, d}; - bool ok = syclex::is_source_kernel_bundle_supported( - ctx.get_backend(), syclex::source_language::opencl); + bool ok = + q.get_device().ext_oneapi_can_compile(syclex::source_language::opencl); if (!ok) { return; } diff --git a/sycl/test-e2e/KernelCompiler/kernel_compiler_spirv.cpp b/sycl/test-e2e/KernelCompiler/kernel_compiler_spirv.cpp index c0e8a7dda85a..38567fe6ee0b 100644 --- a/sycl/test-e2e/KernelCompiler/kernel_compiler_spirv.cpp +++ b/sycl/test-e2e/KernelCompiler/kernel_compiler_spirv.cpp @@ -175,9 +175,7 @@ void testKernelsFromSpvFile(std::string kernels_file, return bundle.ext_oneapi_get_kernel(name); }; - sycl::device d; - sycl::context ctx{d}; - sycl::queue q{ctx, d}; + sycl::queue q; auto bundle = loadKernelsFromFile(q, kernels_file); // Test simple kernel. diff --git a/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl.cpp b/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl.cpp new file mode 100644 index 000000000000..fc294f49fad4 --- /dev/null +++ b/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl.cpp @@ -0,0 +1,276 @@ +//==- kernel_compiler_sycl.cpp --- kernel_compiler extension tests -------==// +// +// 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 +// +//===----------------------------------------------------------------------===// + +// REQUIRES: (opencl || level_zero) +// UNSUPPORTED: accelerator + +// RUN: %{build} -o %t.out +// RUN: %{run} %t.out + +#include +#include + +auto constexpr AddEmH = R"===( + int AddEm(int a, int b){ + return a + b + 5; + } +)==="; + +// TODO: remove SYCL_EXTERNAL once it is no longer needed. +auto constexpr SYCLSource = R"===( +#include +#include "AddEm.h" + +// use extern "C" to avoid name mangling +extern "C" SYCL_EXTERNAL SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((sycl::ext::oneapi::experimental::nd_range_kernel<1>)) +void ff_cp(int *ptr) { + + // intentionally using deprecated routine, as opposed to this_work_item::get_nd_item<1>() + sycl::nd_item<1> Item = sycl::ext::oneapi::experimental::this_nd_item<1>(); + + sycl::id<1> GId = Item.get_global_id(); + ptr[GId.get(0)] = AddEm(GId.get(0), 37); +} + +// this name will be mangled +template +SYCL_EXTERNAL SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((sycl::ext::oneapi::experimental::nd_range_kernel<1>)) +void ff_templated(T *ptr) { + + sycl::nd_item<1> Item = sycl::ext::oneapi::this_work_item::get_nd_item<1>(); + + sycl::id<1> GId = Item.get_global_id(); + ptr[GId.get(0)] = GId.get(0) + 39; +} +)==="; + +auto constexpr BadSource = R"===( +#include + +extern "C" SYCL_EXTERNAL SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((sycl::ext::oneapi::experimental::nd_range_kernel<1>)) +void ff_cp(int *ptr) { + + sycl::nd_item<1> Item = sycl::ext::oneapi::this_work_item::get_nd_item<1>(); + + sycl::id<1> GId = Item.get_global_id() + no semi colon !! + ptr[GId.get(0)] = GId.get(0) + 41; +} +)==="; + +auto constexpr ESIMDSource = R"===( +#include +#include + +using namespace sycl::ext::intel::esimd; + +constexpr int VL = 16; + +extern "C" SYCL_EXTERNAL SYCL_ESIMD_KERNEL SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((sycl::ext::oneapi::experimental::nd_range_kernel<1>)) +void vector_add_esimd(float *A, float *B, float *C) { + sycl::nd_item<1> item = sycl::ext::oneapi::this_work_item::get_nd_item<1>(); + unsigned int i = item.get_global_id(0); + unsigned int offset = i * VL ; + + simd va(A + offset); + simd vb(B + offset); + simd vc = va + vb; + vc.copy_to(C + offset); +} +)==="; + +void test_1(sycl::queue &Queue, sycl::kernel &Kernel, int seed) { + constexpr int Range = 10; + int *usmPtr = sycl::malloc_shared(Range, Queue); + int start = 3; + + sycl::nd_range<1> R1{{Range}, {1}}; + + bool Passa = true; + + memset(usmPtr, 0, Range * sizeof(int)); + Queue.submit([&](sycl::handler &Handler) { + Handler.set_arg(0, usmPtr); + Handler.parallel_for(R1, Kernel); + }); + Queue.wait(); + + for (int i = 0; i < Range; i++) { + std::cout << usmPtr[i] << " "; + assert(usmPtr[i] = i + seed); + } + std::cout << std::endl; + + sycl::free(usmPtr, Queue); +} + +void test_build_and_run() { + namespace syclex = sycl::ext::oneapi::experimental; + using source_kb = sycl::kernel_bundle; + using exe_kb = sycl::kernel_bundle; + + sycl::queue q; + sycl::context ctx = q.get_context(); + + bool ok = + q.get_device().ext_oneapi_can_compile(syclex::source_language::sycl); + if (!ok) { + std::cout << "Apparently this device does not support SYCL source " + "kernel bundle extension: " + << q.get_device().get_info() + << std::endl; + return; + } + + // Create from source. + source_kb kbSrc = syclex::create_kernel_bundle_from_source( + ctx, syclex::source_language::sycl, SYCLSource, + syclex::properties{syclex::include_files{"AddEm.h", AddEmH}}); + + // Double check kernel_bundle.get_source() / get_backend(). + sycl::context ctxRes = kbSrc.get_context(); + assert(ctxRes == ctx); + sycl::backend beRes = kbSrc.get_backend(); + assert(beRes == ctx.get_backend()); + + // Compilation of empty prop list, no devices. + exe_kb kbExe1 = syclex::build(kbSrc); + + // Compilation with props and devices + std::string log; + std::vector flags{"-g", "-fno-fast-math"}; + std::vector devs = kbSrc.get_devices(); + exe_kb kbExe2 = syclex::build( + kbSrc, devs, + syclex::properties{// syclex::build_options{flags}, + syclex::save_log{&log}, + syclex::registered_kernel_names{"ff_templated"}}); + assert(log.find("warning: 'this_nd_item<1>' is deprecated") != + std::string::npos); + + // clang-format off + + // 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"); + + // 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_"); + + // clang-format on + + // Test the kernels. + test_1(q, k, 37 + 5); // AddEm will add 5 more. + test_1(q, k2, 39); +} + +void test_error() { + namespace syclex = sycl::ext::oneapi::experimental; + using source_kb = sycl::kernel_bundle; + using exe_kb = sycl::kernel_bundle; + + sycl::queue q; + sycl::context ctx = q.get_context(); + + bool ok = + q.get_device().ext_oneapi_can_compile(syclex::source_language::sycl); + if (!ok) { + return; + } + + source_kb kbSrc = syclex::create_kernel_bundle_from_source( + ctx, syclex::source_language::sycl, BadSource); + try { + exe_kb kbExe = syclex::build(kbSrc); + assert(false && "we should not be here"); + } catch (sycl::exception &e) { + // yas! + assert(e.code() == sycl::errc::build); + assert(std::string(e.what()).find( + "error: expected ';' at end of declaration") != + std::string::npos); + } +} + +void test_esimd() { + namespace syclex = sycl::ext::oneapi::experimental; + using source_kb = sycl::kernel_bundle; + using exe_kb = sycl::kernel_bundle; + + sycl::queue q; + sycl::context ctx = q.get_context(); + + if (!q.get_device().has(sycl::aspect::ext_intel_esimd)) { + std::cout << "Device does not support ESIMD" << std::endl; + return; + } + + bool ok = + q.get_device().ext_oneapi_can_compile(syclex::source_language::sycl); + if (!ok) { + return; + } + + std::string log; + + source_kb kbSrc = syclex::create_kernel_bundle_from_source( + ctx, syclex::source_language::sycl, ESIMDSource); + exe_kb kbExe = + syclex::build(kbSrc, syclex::properties{syclex::save_log{&log}}); + + // extern "C" was used, so the name "vector_add_esimd" is not mangled and can + // be used directly. + sycl::kernel k = kbExe.ext_oneapi_get_kernel("vector_add_esimd"); + + // Now test it. + constexpr int VL = 16; // this constant also in ESIMDSource string. + constexpr int size = VL * 16; + + float *A = sycl::malloc_shared(size, q); + float *B = sycl::malloc_shared(size, q); + float *C = sycl::malloc_shared(size, q); + for (size_t i = 0; i < size; i++) { + A[i] = float(1); + B[i] = float(2); + C[i] = 0.0f; + } + sycl::range<1> GlobalRange{size / VL}; + sycl::range<1> LocalRange{1}; + sycl::nd_range<1> NDRange{GlobalRange, LocalRange}; + + q.submit([&](sycl::handler &h) { + h.set_arg(0, A); + h.set_arg(1, B); + h.set_arg(2, C); + h.parallel_for(NDRange, k); + }).wait(); + + // Check. + for (size_t i = 0; i < size; i++) { + assert(C[i] == 3.0f); + } + + sycl::free(A, q); + sycl::free(B, q); + sycl::free(C, q); +} + +int main() { + +#ifdef SYCL_EXT_ONEAPI_KERNEL_COMPILER + test_build_and_run(); + test_error(); + test_esimd(); +#else + static_assert(false, "Kernel Compiler feature test macro undefined"); +#endif + return 0; +} diff --git a/sycl/test-e2e/no_sycl_hpp_in_e2e_tests.cpp b/sycl/test-e2e/no_sycl_hpp_in_e2e_tests.cpp index 7b2aa6b3b588..74f8e1aed8af 100644 --- a/sycl/test-e2e/no_sycl_hpp_in_e2e_tests.cpp +++ b/sycl/test-e2e/no_sycl_hpp_in_e2e_tests.cpp @@ -7,7 +7,7 @@ // CHECK-DAG: no_sycl_hpp_in_e2e_tests.cpp // CHECK-DAG: lit.cfg.py // -// CHECK-NUM-MATCHES: 3 +// CHECK-NUM-MATCHES: 4 // // This test verifies that `` isn't used in E2E tests. Instead, // fine-grained includes should used, see diff --git a/sycl/test/abi/sycl_symbols_linux.dump b/sycl/test/abi/sycl_symbols_linux.dump index 99fb95d92fa7..4a7d0dcdcb0e 100644 --- a/sycl/test/abi/sycl_symbols_linux.dump +++ b/sycl/test/abi/sycl_symbols_linux.dump @@ -2988,17 +2988,8 @@ _ZN4sycl3_V13ext5intel12experimental15online_compilerILNS3_15source_languageE0EE _ZN4sycl3_V13ext5intel12experimental15online_compilerILNS3_15source_languageE1EE7compileIJSt6vectorINSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEESaISE_EEEEES8_IhSaIhEERKSE_DpRKT_ _ZN4sycl3_V13ext5intel12experimental9pipe_base13get_pipe_nameB5cxx11EPKv _ZN4sycl3_V13ext5intel12experimental9pipe_base17wait_non_blockingERKNS0_5eventE -_ZN4sycl3_V13ext6oneapi12experimental10mem_adviseENS0_5queueEPvmiRKNS0_6detail13code_locationE _ZN4sycl3_V13ext6oneapi10level_zero6detail11make_deviceERKNS0_8platformEm -_ZN4sycl3_V13ext6oneapi12experimental12physical_memC1ERKNS0_6deviceERKNS0_7contextEm -_ZN4sycl3_V13ext6oneapi12experimental12physical_memC2ERKNS0_6deviceERKNS0_7contextEm -_ZN4sycl3_V13ext6oneapi12experimental15get_access_modeEPKvmRKNS0_7contextE -_ZN4sycl3_V13ext6oneapi12experimental15set_access_modeEPKvmNS3_19address_access_modeERKNS0_7contextE -_ZN4sycl3_V13ext6oneapi12experimental16free_virtual_memEmmRKNS0_7contextE -_ZN4sycl3_V13ext6oneapi12experimental19get_mem_granularityERKNS0_6deviceERKNS0_7contextENS3_16granularity_modeE -_ZN4sycl3_V13ext6oneapi12experimental19get_mem_granularityERKNS0_7contextENS3_16granularity_modeE -_ZN4sycl3_V13ext6oneapi12experimental19reserve_virtual_memEmmRKNS0_7contextE -_ZN4sycl3_V13ext6oneapi12experimental5unmapEPKvmRKNS0_7contextE +_ZN4sycl3_V13ext6oneapi12experimental10mem_adviseENS0_5queueEPvmiRKNS0_6detail13code_locationE _ZN4sycl3_V13ext6oneapi12experimental12create_imageENS3_16image_mem_handleERKNS3_16image_descriptorERKNS0_5queueE _ZN4sycl3_V13ext6oneapi12experimental12create_imageENS3_16image_mem_handleERKNS3_16image_descriptorERKNS0_6deviceERKNS0_7contextE _ZN4sycl3_V13ext6oneapi12experimental12create_imageENS3_16image_mem_handleERKNS3_22bindless_image_samplerERKNS3_16image_descriptorERKNS0_5queueE @@ -3009,6 +3000,8 @@ _ZN4sycl3_V13ext6oneapi12experimental12create_imageERNS3_9image_memERKNS3_16imag _ZN4sycl3_V13ext6oneapi12experimental12create_imageERNS3_9image_memERKNS3_16image_descriptorERKNS0_6deviceERKNS0_7contextE _ZN4sycl3_V13ext6oneapi12experimental12create_imageERNS3_9image_memERKNS3_22bindless_image_samplerERKNS3_16image_descriptorERKNS0_5queueE _ZN4sycl3_V13ext6oneapi12experimental12create_imageERNS3_9image_memERKNS3_22bindless_image_samplerERKNS3_16image_descriptorERKNS0_6deviceERKNS0_7contextE +_ZN4sycl3_V13ext6oneapi12experimental12physical_memC1ERKNS0_6deviceERKNS0_7contextEm +_ZN4sycl3_V13ext6oneapi12experimental12physical_memC2ERKNS0_6deviceERKNS0_7contextEm _ZN4sycl3_V13ext6oneapi12experimental14free_image_memENS3_16image_mem_handleENS3_10image_typeERKNS0_5queueE _ZN4sycl3_V13ext6oneapi12experimental14free_image_memENS3_16image_mem_handleENS3_10image_typeERKNS0_6deviceERKNS0_7contextE _ZN4sycl3_V13ext6oneapi12experimental14free_image_memENS3_16image_mem_handleERKNS0_5queueE @@ -3017,10 +3010,16 @@ _ZN4sycl3_V13ext6oneapi12experimental15alloc_image_memERKNS3_16image_descriptorE _ZN4sycl3_V13ext6oneapi12experimental15alloc_image_memERKNS3_16image_descriptorERKNS0_6deviceERKNS0_7contextE _ZN4sycl3_V13ext6oneapi12experimental15free_mipmap_memENS3_16image_mem_handleERKNS0_5queueE _ZN4sycl3_V13ext6oneapi12experimental15free_mipmap_memENS3_16image_mem_handleERKNS0_6deviceERKNS0_7contextE +_ZN4sycl3_V13ext6oneapi12experimental15get_access_modeEPKvmRKNS0_7contextE _ZN4sycl3_V13ext6oneapi12experimental15get_image_rangeENS3_16image_mem_handleERKNS0_5queueE _ZN4sycl3_V13ext6oneapi12experimental15get_image_rangeENS3_16image_mem_handleERKNS0_6deviceERKNS0_7contextE +_ZN4sycl3_V13ext6oneapi12experimental15set_access_modeEPKvmNS3_19address_access_modeERKNS0_7contextE _ZN4sycl3_V13ext6oneapi12experimental16alloc_mipmap_memERKNS3_16image_descriptorERKNS0_5queueE _ZN4sycl3_V13ext6oneapi12experimental16alloc_mipmap_memERKNS3_16image_descriptorERKNS0_6deviceERKNS0_7contextE +_ZN4sycl3_V13ext6oneapi12experimental16free_virtual_memEmmRKNS0_7contextE +_ZN4sycl3_V13ext6oneapi12experimental19get_mem_granularityERKNS0_6deviceERKNS0_7contextENS3_16granularity_modeE +_ZN4sycl3_V13ext6oneapi12experimental19get_mem_granularityERKNS0_7contextENS3_16granularity_modeE +_ZN4sycl3_V13ext6oneapi12experimental19reserve_virtual_memEmmRKNS0_7contextE _ZN4sycl3_V13ext6oneapi12experimental20destroy_image_handleERNS3_20sampled_image_handleERKNS0_5queueE _ZN4sycl3_V13ext6oneapi12experimental20destroy_image_handleERNS3_20sampled_image_handleERKNS0_6deviceERKNS0_7contextE _ZN4sycl3_V13ext6oneapi12experimental20destroy_image_handleERNS3_22unsampled_image_handleERKNS0_5queueE @@ -3060,9 +3059,6 @@ _ZN4sycl3_V13ext6oneapi12experimental25map_external_memory_arrayENS3_18interop_m _ZN4sycl3_V13ext6oneapi12experimental25map_external_memory_arrayENS3_18interop_mem_handleERKNS3_16image_descriptorERKNS0_6deviceERKNS0_7contextE _ZN4sycl3_V13ext6oneapi12experimental26destroy_external_semaphoreENS3_24interop_semaphore_handleERKNS0_5queueE _ZN4sycl3_V13ext6oneapi12experimental26destroy_external_semaphoreENS3_24interop_semaphore_handleERKNS0_6deviceERKNS0_7contextE -_ZN4sycl3_V13ext6oneapi12experimental32create_kernel_bundle_from_sourceERKNS0_7contextENS3_15source_languageERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEE -_ZN4sycl3_V13ext6oneapi12experimental32create_kernel_bundle_from_sourceERKNS0_7contextENS3_15source_languageERKSt6vectorISt4byteSaIS9_EE -_ZN4sycl3_V13ext6oneapi12experimental33is_source_kernel_bundle_supportedENS0_7backendENS3_15source_languageE _ZN4sycl3_V13ext6oneapi12experimental4node12update_rangeILi1EEEvNS0_5rangeIXT_EEE _ZN4sycl3_V13ext6oneapi12experimental4node12update_rangeILi2EEEvNS0_5rangeIXT_EEE _ZN4sycl3_V13ext6oneapi12experimental4node12update_rangeILi3EEEvNS0_5rangeIXT_EEE @@ -3070,11 +3066,12 @@ _ZN4sycl3_V13ext6oneapi12experimental4node15update_nd_rangeILi1EEEvNS0_8nd_range _ZN4sycl3_V13ext6oneapi12experimental4node15update_nd_rangeILi2EEEvNS0_8nd_rangeIXT_EEE _ZN4sycl3_V13ext6oneapi12experimental4node15update_nd_rangeILi3EEEvNS0_8nd_rangeIXT_EEE _ZN4sycl3_V13ext6oneapi12experimental4node19get_node_from_eventENS0_5eventE +_ZN4sycl3_V13ext6oneapi12experimental5unmapEPKvmRKNS0_7contextE _ZN4sycl3_V13ext6oneapi12experimental6detail14image_mem_implC1ERKNS3_16image_descriptorERKNS0_6deviceERKNS0_7contextE _ZN4sycl3_V13ext6oneapi12experimental6detail14image_mem_implC2ERKNS3_16image_descriptorERKNS0_6deviceERKNS0_7contextE _ZN4sycl3_V13ext6oneapi12experimental6detail14image_mem_implD1Ev _ZN4sycl3_V13ext6oneapi12experimental6detail14image_mem_implD2Ev -_ZN4sycl3_V13ext6oneapi12experimental6detail17build_from_sourceERNS0_13kernel_bundleILNS0_12bundle_stateE3EEERKSt6vectorINS0_6deviceESaISA_EERKS9_INSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEESaISK_EEPSK_ +_ZN4sycl3_V13ext6oneapi12experimental6detail17build_from_sourceERNS0_13kernel_bundleILNS0_12bundle_stateE3EEERKSt6vectorINS0_6deviceESaISA_EERKS9_INSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEESaISK_EEPSK_SO_ _ZN4sycl3_V13ext6oneapi12experimental6detail22dynamic_parameter_base11updateValueEPKvm _ZN4sycl3_V13ext6oneapi12experimental6detail22dynamic_parameter_base14updateAccessorEPKNS0_6detail16AccessorBaseHostE _ZN4sycl3_V13ext6oneapi12experimental6detail22dynamic_parameter_baseC1ENS3_13command_graphILNS3_11graph_stateE0EEEmPKv @@ -3098,6 +3095,9 @@ _ZN4sycl3_V13ext6oneapi12experimental6detail24modifiable_command_graphC1ERKNS0_5 _ZN4sycl3_V13ext6oneapi12experimental6detail24modifiable_command_graphC1ERKNS0_7contextERKNS0_6deviceERKNS0_13property_listE _ZN4sycl3_V13ext6oneapi12experimental6detail24modifiable_command_graphC2ERKNS0_5queueERKNS0_13property_listE _ZN4sycl3_V13ext6oneapi12experimental6detail24modifiable_command_graphC2ERKNS0_7contextERKNS0_6deviceERKNS0_13property_listE +_ZN4sycl3_V13ext6oneapi12experimental6detail30make_kernel_bundle_from_sourceERKNS0_7contextENS3_15source_languageERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEESt6vectorISt4pairISE_SE_ESaISJ_EE +_ZN4sycl3_V13ext6oneapi12experimental6detail30make_kernel_bundle_from_sourceERKNS0_7contextENS3_15source_languageERKSt6vectorISt4byteSaISA_EES9_ISt4pairINSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEESL_ESaISM_EE +_ZN4sycl3_V13ext6oneapi12experimental6detail33is_source_kernel_bundle_supportedENS0_7backendENS3_15source_languageE _ZN4sycl3_V13ext6oneapi12experimental6memcpyENS0_5queueEPvPKvmRKNS0_6detail13code_locationE _ZN4sycl3_V13ext6oneapi12experimental6memsetENS0_5queueEPvimRKNS0_6detail13code_locationE _ZN4sycl3_V13ext6oneapi12experimental9image_memC1ERKNS3_16image_descriptorERKNS0_5queueE diff --git a/sycl/test/abi/sycl_symbols_windows.dump b/sycl/test/abi/sycl_symbols_windows.dump index 9b80d2eb69c8..00d45e891377 100644 --- a/sycl/test/abi/sycl_symbols_windows.dump +++ b/sycl/test/abi/sycl_symbols_windows.dump @@ -569,10 +569,10 @@ ??0half@host_half_impl@detail@_V1@sycl@@QEAA@AEBM@Z ??0half@host_half_impl@detail@_V1@sycl@@QEAA@G@Z ??0handler@_V1@sycl@@AEAA@V?$shared_ptr@Vgraph_impl@detail@experimental@oneapi@ext@_V1@sycl@@@std@@@Z -??0handler@_V1@sycl@@AEAA@V?$shared_ptr@Vqueue_impl@detail@_V1@sycl@@@std@@00_N@Z ??0handler@_V1@sycl@@AEAA@V?$shared_ptr@Vqueue_impl@detail@_V1@sycl@@@std@@00_N1@Z -??0handler@_V1@sycl@@AEAA@V?$shared_ptr@Vqueue_impl@detail@_V1@sycl@@@std@@_N@Z +??0handler@_V1@sycl@@AEAA@V?$shared_ptr@Vqueue_impl@detail@_V1@sycl@@@std@@00_N@Z ??0handler@_V1@sycl@@AEAA@V?$shared_ptr@Vqueue_impl@detail@_V1@sycl@@@std@@_N1@Z +??0handler@_V1@sycl@@AEAA@V?$shared_ptr@Vqueue_impl@detail@_V1@sycl@@@std@@_N@Z ??0host_selector@_V1@sycl@@QEAA@$$QEAV012@@Z ??0host_selector@_V1@sycl@@QEAA@AEBV012@@Z ??0host_selector@_V1@sycl@@QEAA@XZ @@ -609,10 +609,6 @@ ??0kernel_id@_V1@sycl@@AEAA@PEBD@Z ??0kernel_id@_V1@sycl@@QEAA@$$QEAV012@@Z ??0kernel_id@_V1@sycl@@QEAA@AEBV012@@Z -??0physical_mem@experimental@oneapi@ext@_V1@sycl@@QEAA@AEBV012345@@Z -??0physical_mem@experimental@oneapi@ext@_V1@sycl@@QEAA@$$QEAV012345@@Z -??0physical_mem@experimental@oneapi@ext@_V1@sycl@@QEAA@AEBVqueue@45@_K@Z -??0physical_mem@experimental@oneapi@ext@_V1@sycl@@QEAA@AEBVdevice@45@AEBVcontext@45@_K@Z ??0modifiable_command_graph@detail@experimental@oneapi@ext@_V1@sycl@@IEAA@AEBV?$shared_ptr@Vgraph_impl@detail@experimental@oneapi@ext@_V1@sycl@@@std@@@Z ??0modifiable_command_graph@detail@experimental@oneapi@ext@_V1@sycl@@QEAA@$$QEAV0123456@@Z ??0modifiable_command_graph@detail@experimental@oneapi@ext@_V1@sycl@@QEAA@AEBV0123456@@Z @@ -621,6 +617,10 @@ ??0node@experimental@oneapi@ext@_V1@sycl@@AEAA@AEBV?$shared_ptr@Vnode_impl@detail@experimental@oneapi@ext@_V1@sycl@@@std@@@Z ??0node@experimental@oneapi@ext@_V1@sycl@@QEAA@$$QEAV012345@@Z ??0node@experimental@oneapi@ext@_V1@sycl@@QEAA@AEBV012345@@Z +??0physical_mem@experimental@oneapi@ext@_V1@sycl@@QEAA@$$QEAV012345@@Z +??0physical_mem@experimental@oneapi@ext@_V1@sycl@@QEAA@AEBV012345@@Z +??0physical_mem@experimental@oneapi@ext@_V1@sycl@@QEAA@AEBVdevice@45@AEBVcontext@45@_K@Z +??0physical_mem@experimental@oneapi@ext@_V1@sycl@@QEAA@AEBVqueue@45@_K@Z ??0platform@_V1@sycl@@AEAA@AEBVdevice@12@@Z ??0platform@_V1@sycl@@AEAA@V?$shared_ptr@Vplatform_impl@detail@_V1@sycl@@@std@@@Z ??0platform@_V1@sycl@@QEAA@$$QEAV012@@Z @@ -683,9 +683,9 @@ ??1kernel@_V1@sycl@@QEAA@XZ ??1kernel_bundle_plain@detail@_V1@sycl@@QEAA@XZ ??1kernel_id@_V1@sycl@@QEAA@XZ -??1physical_mem@experimental@oneapi@ext@_V1@sycl@@QEAA@XZ ??1modifiable_command_graph@detail@experimental@oneapi@ext@_V1@sycl@@QEAA@XZ ??1node@experimental@oneapi@ext@_V1@sycl@@QEAA@XZ +??1physical_mem@experimental@oneapi@ext@_V1@sycl@@QEAA@XZ ??1platform@_V1@sycl@@QEAA@XZ ??1queue@_V1@sycl@@QEAA@XZ ??1sampler@_V1@sycl@@QEAA@XZ @@ -768,12 +768,12 @@ ??4kernel_bundle_plain@detail@_V1@sycl@@QEAAAEAV0123@AEBV0123@@Z ??4kernel_id@_V1@sycl@@QEAAAEAV012@$$QEAV012@@Z ??4kernel_id@_V1@sycl@@QEAAAEAV012@AEBV012@@Z -??4physical_mem@experimental@oneapi@ext@_V1@sycl@@QEAAAEAV012345@$$QEAV012345@@Z -??4physical_mem@experimental@oneapi@ext@_V1@sycl@@QEAAAEAV012345@AEBV012345@@Z ??4modifiable_command_graph@detail@experimental@oneapi@ext@_V1@sycl@@QEAAAEAV0123456@$$QEAV0123456@@Z ??4modifiable_command_graph@detail@experimental@oneapi@ext@_V1@sycl@@QEAAAEAV0123456@AEBV0123456@@Z ??4node@experimental@oneapi@ext@_V1@sycl@@QEAAAEAV012345@$$QEAV012345@@Z ??4node@experimental@oneapi@ext@_V1@sycl@@QEAAAEAV012345@AEBV012345@@Z +??4physical_mem@experimental@oneapi@ext@_V1@sycl@@QEAAAEAV012345@$$QEAV012345@@Z +??4physical_mem@experimental@oneapi@ext@_V1@sycl@@QEAAAEAV012345@AEBV012345@@Z ??4platform@_V1@sycl@@QEAAAEAV012@$$QEAV012@@Z ??4platform@_V1@sycl@@QEAAAEAV012@AEBV012@@Z ??4queue@_V1@sycl@@QEAAAEAV012@$$QEAV012@@Z @@ -3933,7 +3933,7 @@ ?begin@kernel_bundle_plain@detail@_V1@sycl@@IEBAPEBVdevice_image_plain@234@XZ ?begin_recording@modifiable_command_graph@detail@experimental@oneapi@ext@_V1@sycl@@QEAAXAEAVqueue@67@AEBVproperty_list@67@@Z ?begin_recording@modifiable_command_graph@detail@experimental@oneapi@ext@_V1@sycl@@QEAAXAEBV?$vector@Vqueue@_V1@sycl@@V?$allocator@Vqueue@_V1@sycl@@@std@@@std@@AEBVproperty_list@67@@Z -?build_from_source@detail@experimental@oneapi@ext@_V1@sycl@@YA?AV?$kernel_bundle@$01@56@AEAV?$kernel_bundle@$02@56@AEBV?$vector@Vdevice@_V1@sycl@@V?$allocator@Vdevice@_V1@sycl@@@std@@@std@@AEBV?$vector@V?$basic_string@DU?$char_traits@D@std@@V?$allocator@D@2@@std@@V?$allocator@V?$basic_string@DU?$char_traits@D@std@@V?$allocator@D@2@@std@@@2@@std@@PEAV?$basic_string@DU?$char_traits@D@std@@V?$allocator@D@2@@std@@@Z +?build_from_source@detail@experimental@oneapi@ext@_V1@sycl@@YA?AV?$kernel_bundle@$01@56@AEAV?$kernel_bundle@$02@56@AEBV?$vector@Vdevice@_V1@sycl@@V?$allocator@Vdevice@_V1@sycl@@@std@@@std@@AEBV?$vector@V?$basic_string@DU?$char_traits@D@std@@V?$allocator@D@2@@std@@V?$allocator@V?$basic_string@DU?$char_traits@D@std@@V?$allocator@D@2@@std@@@2@@std@@PEAV?$basic_string@DU?$char_traits@D@std@@V?$allocator@D@2@@std@@2@Z ?build_impl@detail@_V1@sycl@@YA?AV?$shared_ptr@Vkernel_bundle_impl@detail@_V1@sycl@@@std@@AEBV?$kernel_bundle@$0A@@23@AEBV?$vector@Vdevice@_V1@sycl@@V?$allocator@Vdevice@_V1@sycl@@@std@@@5@AEBVproperty_list@23@@Z ?cancel_fusion@fusion_wrapper@experimental@codeplay@ext@_V1@sycl@@QEAAXXZ ?category@exception@_V1@sycl@@QEBAAEBVerror_category@std@@XZ @@ -3956,8 +3956,6 @@ ?create_image@experimental@oneapi@ext@_V1@sycl@@YA?AUunsampled_image_handle@12345@AEAVimage_mem@12345@AEBUimage_descriptor@12345@AEBVqueue@45@@Z ?create_image@experimental@oneapi@ext@_V1@sycl@@YA?AUunsampled_image_handle@12345@Uimage_mem_handle@12345@AEBUimage_descriptor@12345@AEBVdevice@45@AEBVcontext@45@@Z ?create_image@experimental@oneapi@ext@_V1@sycl@@YA?AUunsampled_image_handle@12345@Uimage_mem_handle@12345@AEBUimage_descriptor@12345@AEBVqueue@45@@Z -?create_kernel_bundle_from_source@experimental@oneapi@ext@_V1@sycl@@YA?AV?$kernel_bundle@$02@45@AEBVcontext@45@W4source_language@12345@AEBV?$basic_string@DU?$char_traits@D@std@@V?$allocator@D@2@@std@@@Z -?create_kernel_bundle_from_source@experimental@oneapi@ext@_V1@sycl@@YA?AV?$kernel_bundle@$02@45@AEBVcontext@45@W4source_language@12345@AEBV?$vector@W4byte@std@@V?$allocator@W4byte@std@@@2@@std@@@Z ?default_selector_v@_V1@sycl@@YAHAEBVdevice@12@@Z ?deleteAccProps@buffer_plain@detail@_V1@sycl@@IEAAXAEBW4PropWithDataKind@234@@Z ?depends_on@handler@_V1@sycl@@IEAAXAEBV?$shared_ptr@Vevent_impl@detail@_V1@sycl@@@std@@@Z @@ -4088,14 +4086,13 @@ ?find_device_intersection@detail@_V1@sycl@@YA?AV?$vector@Vdevice@_V1@sycl@@V?$allocator@Vdevice@_V1@sycl@@@std@@@std@@AEBV?$vector@V?$kernel_bundle@$00@_V1@sycl@@V?$allocator@V?$kernel_bundle@$00@_V1@sycl@@@std@@@5@@Z ?free@_V1@sycl@@YAXPEAXAEBVcontext@12@AEBUcode_location@detail@12@@Z ?free@_V1@sycl@@YAXPEAXAEBVqueue@12@AEBUcode_location@detail@12@@Z -?free_virtual_mem@experimental@oneapi@ext@_V1@sycl@@YAX_K0AEBVcontext@45@@Z +?free_image_mem@experimental@oneapi@ext@_V1@sycl@@YAXUimage_mem_handle@12345@AEBVdevice@45@AEBVcontext@45@@Z ?free_image_mem@experimental@oneapi@ext@_V1@sycl@@YAXUimage_mem_handle@12345@AEBVqueue@45@@Z ?free_image_mem@experimental@oneapi@ext@_V1@sycl@@YAXUimage_mem_handle@12345@W4image_type@12345@AEBVdevice@45@AEBVcontext@45@@Z ?free_image_mem@experimental@oneapi@ext@_V1@sycl@@YAXUimage_mem_handle@12345@W4image_type@12345@AEBVqueue@45@@Z -?free_image_mem@experimental@oneapi@ext@_V1@sycl@@YAXUimage_mem_handle@12345@AEBVdevice@45@AEBVcontext@45@@Z ?free_mipmap_mem@experimental@oneapi@ext@_V1@sycl@@YAXUimage_mem_handle@12345@AEBVdevice@45@AEBVcontext@45@@Z ?free_mipmap_mem@experimental@oneapi@ext@_V1@sycl@@YAXUimage_mem_handle@12345@AEBVqueue@45@@Z -?free_mipmap_mem@experimental@oneapi@ext@_V1@sycl@@YAXUimage_mem_handle@12345@AEBVdevice@45@AEBVcontext@45@@Z +?free_virtual_mem@experimental@oneapi@ext@_V1@sycl@@YAX_K0AEBVcontext@45@@Z ?frexp_impl@detail@_V1@sycl@@YA?AVhalf@half_impl@123@V45123@PEAH@Z ?frexp_impl@detail@_V1@sycl@@YAMMPEAH@Z ?frexp_impl@detail@_V1@sycl@@YANNPEAH@Z @@ -4210,8 +4207,8 @@ ?get_coordinate_normalization_mode@sampler@_V1@sycl@@QEBA?AW4coordinate_normalization_mode@23@XZ ?get_count@image_plain@detail@_V1@sycl@@IEBA_KXZ ?get_descriptor@image_mem@experimental@oneapi@ext@_V1@sycl@@QEBAAEBUimage_descriptor@23456@XZ -?get_device@physical_mem@experimental@oneapi@ext@_V1@sycl@@QEBA?AVdevice@56@XZ ?get_device@image_mem@experimental@oneapi@ext@_V1@sycl@@QEBA?AVdevice@56@XZ +?get_device@physical_mem@experimental@oneapi@ext@_V1@sycl@@QEBA?AVdevice@56@XZ ?get_device@queue@_V1@sycl@@QEBA?AVdevice@23@XZ ?get_devices@context@_V1@sycl@@QEBA?AV?$vector@Vdevice@_V1@sycl@@V?$allocator@Vdevice@_V1@sycl@@@std@@@std@@XZ ?get_devices@device@_V1@sycl@@SA?AV?$vector@Vdevice@_V1@sycl@@V?$allocator@Vdevice@_V1@sycl@@@std@@@std@@W4device_type@info@23@@Z @@ -4310,7 +4307,7 @@ ?is_host@queue@_V1@sycl@@QEBA_NXZ ?is_in_fusion_mode@fusion_wrapper@experimental@codeplay@ext@_V1@sycl@@QEBA_NXZ ?is_in_order@queue@_V1@sycl@@QEBA_NXZ -?is_source_kernel_bundle_supported@experimental@oneapi@ext@_V1@sycl@@YA_NW4backend@45@W4source_language@12345@@Z +?is_source_kernel_bundle_supported@detail@experimental@oneapi@ext@_V1@sycl@@YA_NW4backend@56@W4source_language@23456@@Z ?is_specialization_constant_set@kernel_bundle_plain@detail@_V1@sycl@@IEBA_NPEBD@Z ?join_impl@detail@_V1@sycl@@YA?AV?$shared_ptr@Vkernel_bundle_impl@detail@_V1@sycl@@@std@@AEBV?$vector@V?$shared_ptr@Vkernel_bundle_impl@detail@_V1@sycl@@@std@@V?$allocator@V?$shared_ptr@Vkernel_bundle_impl@detail@_V1@sycl@@@std@@@2@@5@W4bundle_state@23@@Z ?lgamma_r_impl@detail@_V1@sycl@@YA?AVhalf@half_impl@123@V45123@PEAH@Z @@ -4329,6 +4326,8 @@ ?make_kernel@detail@_V1@sycl@@YA?AVkernel@23@_KAEBVcontext@23@W4backend@23@@Z ?make_kernel_bundle@detail@_V1@sycl@@YA?AV?$shared_ptr@Vkernel_bundle_impl@detail@_V1@sycl@@@std@@_KAEBVcontext@23@W4bundle_state@23@W4backend@23@@Z ?make_kernel_bundle@detail@_V1@sycl@@YA?AV?$shared_ptr@Vkernel_bundle_impl@detail@_V1@sycl@@@std@@_KAEBVcontext@23@_NW4bundle_state@23@W4backend@23@@Z +?make_kernel_bundle_from_source@detail@experimental@oneapi@ext@_V1@sycl@@YA?AV?$kernel_bundle@$02@56@AEBVcontext@56@W4source_language@23456@AEBV?$basic_string@DU?$char_traits@D@std@@V?$allocator@D@2@@std@@V?$vector@U?$pair@V?$basic_string@DU?$char_traits@D@std@@V?$allocator@D@2@@std@@V12@@std@@V?$allocator@U?$pair@V?$basic_string@DU?$char_traits@D@std@@V?$allocator@D@2@@std@@V12@@std@@@2@@std@@@Z +?make_kernel_bundle_from_source@detail@experimental@oneapi@ext@_V1@sycl@@YA?AV?$kernel_bundle@$02@56@AEBVcontext@56@W4source_language@23456@AEBV?$vector@W4byte@std@@V?$allocator@W4byte@std@@@2@@std@@V?$vector@U?$pair@V?$basic_string@DU?$char_traits@D@std@@V?$allocator@D@2@@std@@V12@@std@@V?$allocator@U?$pair@V?$basic_string@DU?$char_traits@D@std@@V?$allocator@D@2@@std@@V12@@std@@@2@@std@@@Z ?make_platform@detail@_V1@sycl@@YA?AVplatform@23@_KW4backend@23@@Z ?make_queue@detail@_V1@sycl@@YA?AVqueue@23@_KHAEBVcontext@23@PEBVdevice@23@_NAEBVproperty_list@23@AEBV?$function@$$A6AXVexception_list@_V1@sycl@@@Z@std@@W4backend@23@@Z ?malloc@_V1@sycl@@YAPEAX_KAEBVdevice@12@AEBVcontext@12@W4alloc@usm@12@AEBUcode_location@detail@12@@Z