diff --git a/sycl-jit/jit-compiler/lib/rtc/DeviceCompilation.cpp b/sycl-jit/jit-compiler/lib/rtc/DeviceCompilation.cpp index a47e2800ec668..df8812686643c 100644 --- a/sycl-jit/jit-compiler/lib/rtc/DeviceCompilation.cpp +++ b/sycl-jit/jit-compiler/lib/rtc/DeviceCompilation.cpp @@ -657,8 +657,9 @@ jit_compiler::performPostLink(std::unique_ptr Module, /*IROutputOnly=*/false, EmitOnlyKernelsAsEntryPoints); assert(Splitter->hasMoreSplits()); - // TODO: Call `verifyNoCrossModuleDeviceGlobalUsage` if device globals shall - // be processed. + if (auto Err = Splitter->verifyNoCrossModuleDeviceGlobalUsage()) { + return std::move(Err); + } SmallVector DevImgInfoVec; SmallVector> Modules; @@ -698,7 +699,7 @@ jit_compiler::performPostLink(std::unique_ptr Module, GlobalBinImageProps PropReq{ /*EmitKernelParamInfo=*/true, /*EmitProgramMetadata=*/true, /*EmitExportedSymbols=*/true, /*EmitImportedSymbols=*/true, - /*DeviceGlobals=*/false}; + /*DeviceGlobals=*/true}; PropertySetRegistry Properties = computeModuleProperties(MDesc.getModule(), MDesc.entries(), PropReq); // TODO: Manually add `compile_target` property as in diff --git a/sycl/include/sycl/kernel_bundle.hpp b/sycl/include/sycl/kernel_bundle.hpp index c6b9c23b0d2a4..dc3e683d1dabd 100644 --- a/sycl/include/sycl/kernel_bundle.hpp +++ b/sycl/include/sycl/kernel_bundle.hpp @@ -242,6 +242,19 @@ class __SYCL_EXPORT kernel_bundle_plain { ext_oneapi_get_raw_kernel_name(detail::string_view{name}).c_str()}; } + bool ext_oneapi_has_device_global(const std::string &name) { + return ext_oneapi_has_device_global(detail::string_view{name}); + } + + void *ext_oneapi_get_device_global_address(const std::string &name, + const device &dev) { + return ext_oneapi_get_device_global_address(detail::string_view{name}, dev); + } + + size_t ext_oneapi_get_device_global_size(const std::string &name) { + return ext_oneapi_get_device_global_size(detail::string_view{name}); + } + protected: // \returns a kernel object which represents the kernel identified by // kernel_id passed @@ -271,6 +284,11 @@ class __SYCL_EXPORT kernel_bundle_plain { 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); + + bool ext_oneapi_has_device_global(detail::string_view name); + void *ext_oneapi_get_device_global_address(detail::string_view name, + const device &dev); + size_t ext_oneapi_get_device_global_size(detail::string_view name); }; } // namespace detail @@ -501,6 +519,43 @@ class kernel_bundle : public detail::kernel_bundle_plain, return detail::kernel_bundle_plain::ext_oneapi_get_raw_kernel_name(name); } + ///////////////////////// + // ext_oneapi_has_device_global + // only true if kernel_bundle was created from source and has this device + // global + ///////////////////////// + template > + bool ext_oneapi_has_device_global(const std::string &name) { + return detail::kernel_bundle_plain::ext_oneapi_has_device_global(name); + } + + ///////////////////////// + // ext_oneapi_get_device_global_address + // kernel_bundle must be created from source, throws if bundle was not built + // for this device, or device global is either not present or has + // `device_image_scope` property. + // Returns a USM pointer to the variable's initialized storage on the device. + ///////////////////////// + template > + void *ext_oneapi_get_device_global_address(const std::string &name, + const device &dev) { + return detail::kernel_bundle_plain::ext_oneapi_get_device_global_address( + name, dev); + } + + ///////////////////////// + // ext_oneapi_get_device_global_size + // kernel_bundle must be created from source, throws if device global is not + // present. Returns the variable's size in bytes. + ///////////////////////// + template > + size_t ext_oneapi_get_device_global_size(const std::string &name) { + return detail::kernel_bundle_plain::ext_oneapi_get_device_global_size(name); + } + private: kernel_bundle(detail::KernelBundleImplPtr Impl) : kernel_bundle_plain(std::move(Impl)) {} diff --git a/sycl/source/detail/context_impl.cpp b/sycl/source/detail/context_impl.cpp index e527d0a0c46a8..05aace25dface 100644 --- a/sycl/source/detail/context_impl.cpp +++ b/sycl/source/detail/context_impl.cpp @@ -329,6 +329,11 @@ void context_impl::addAssociatedDeviceGlobal(const void *DeviceGlobalPtr) { MAssociatedDeviceGlobals.insert(DeviceGlobalPtr); } +void context_impl::removeAssociatedDeviceGlobal(const void *DeviceGlobalPtr) { + std::lock_guard Lock{MAssociatedDeviceGlobalsMutex}; + MAssociatedDeviceGlobals.erase(DeviceGlobalPtr); +} + void context_impl::addDeviceGlobalInitializer( ur_program_handle_t Program, const std::vector &Devs, const RTDeviceBinaryImage *BinImage) { @@ -407,6 +412,9 @@ std::vector context_impl::initializeDeviceGlobals( // Device global map entry pointers will not die before the end of the // program and the pointers will stay the same, so we do not need // m_DeviceGlobalsMutex here. + // The lifetimes of device global map entries representing globals in + // runtime-compiled code will be tied to the kernel bundle, so the + // assumption holds in that setting as well. for (DeviceGlobalMapEntry *DeviceGlobalEntry : DeviceGlobalEntries) { // Get or allocate the USM memory associated with the device global. DeviceGlobalUSMMem &DeviceGlobalUSM = diff --git a/sycl/source/detail/context_impl.hpp b/sycl/source/detail/context_impl.hpp index 65e1fda0a5a7a..d573d19977ccd 100644 --- a/sycl/source/detail/context_impl.hpp +++ b/sycl/source/detail/context_impl.hpp @@ -203,6 +203,9 @@ class context_impl { /// Adds an associated device global to the tracked associates. void addAssociatedDeviceGlobal(const void *DeviceGlobalPtr); + /// Removes an associated device global from the tracked associates. + void removeAssociatedDeviceGlobal(const void *DeviceGlobalPtr); + /// Adds a device global initializer. void addDeviceGlobalInitializer(ur_program_handle_t Program, const std::vector &Devs, diff --git a/sycl/source/detail/jit_compiler.cpp b/sycl/source/detail/jit_compiler.cpp index 0337b2d98e4b8..93e39db623565 100644 --- a/sycl/source/detail/jit_compiler.cpp +++ b/sycl/source/detail/jit_compiler.cpp @@ -1170,14 +1170,18 @@ sycl_device_binaries jit_compiler::createDeviceBinaries( } for (const auto &FPS : DevImgInfo.Properties) { + bool IsDeviceGlobalsPropSet = + FPS.Name == __SYCL_PROPERTY_SET_SYCL_DEVICE_GLOBALS; PropertySetContainer PropSet{FPS.Name.c_str()}; for (const auto &FPV : FPS.Values) { if (FPV.IsUIntValue) { PropSet.addProperty( PropertyContainer{FPV.Name.c_str(), FPV.UIntValue}); } else { + std::string PrefixedName = + (IsDeviceGlobalsPropSet ? Prefix : "") + FPV.Name.c_str(); PropSet.addProperty(PropertyContainer{ - FPV.Name.c_str(), FPV.Bytes.begin(), FPV.Bytes.size(), + PrefixedName.c_str(), FPV.Bytes.begin(), FPV.Bytes.size(), sycl_property_type::SYCL_PROPERTY_TYPE_BYTE_ARRAY}); } } diff --git a/sycl/source/detail/kernel_bundle_impl.hpp b/sycl/source/detail/kernel_bundle_impl.hpp index 666e6432de8b6..ffb2273c4a75c 100644 --- a/sycl/source/detail/kernel_bundle_impl.hpp +++ b/sycl/source/detail/kernel_bundle_impl.hpp @@ -25,6 +25,7 @@ #include #include #include +#include #include #include "split_string.hpp" @@ -383,6 +384,8 @@ class kernel_bundle_impl { const std::vector &KernelIDs, std::vector &&KernelNames, std::unordered_map &&MangledKernelNames, + std::vector &&DeviceGlobalNames, + std::vector> &&DeviceGlobalAllocations, sycl_device_binaries Binaries, std::string &&Prefix, syclex::source_language Lang) : kernel_bundle_impl(std::move(Ctx), std::move(Devs), KernelIDs, @@ -396,6 +399,8 @@ class kernel_bundle_impl { MIsInterop = true; MKernelNames = std::move(KernelNames); MMangledKernelNames = std::move(MangledKernelNames); + MDeviceGlobalNames = std::move(DeviceGlobalNames); + MDeviceGlobalAllocations = std::move(DeviceGlobalAllocations); MDeviceBinaries = Binaries; MPrefix = std::move(Prefix); MLanguage = Lang; @@ -546,6 +551,12 @@ class kernel_bundle_impl { std::vector KernelIDs; std::vector KernelNames; std::unordered_map MangledKernelNames; + + std::unordered_set DeviceGlobalIDSet; + std::vector DeviceGlobalIDVec; + std::vector DeviceGlobalNames; + std::vector> DeviceGlobalAllocations; + for (const auto &KernelID : PM.getAllSYCLKernelIDs()) { std::string_view KernelName{KernelID.get_name()}; if (KernelName.find(Prefix) == 0) { @@ -563,8 +574,8 @@ class kernel_bundle_impl { } } - // Apply frontend information. for (const auto *RawImg : PM.getRawDeviceImages(KernelIDs)) { + // Mangled names. for (const sycl_device_binary_property &RKProp : RawImg->getRegisteredKernels()) { @@ -574,11 +585,49 @@ class kernel_bundle_impl { reinterpret_cast(BA.begin()), MangledNameLen}; MangledKernelNames.emplace(RKProp->Name, MangledName); } + + // Device globals. + for (const auto &DeviceGlobalProp : RawImg->getDeviceGlobals()) { + std::string_view DeviceGlobalName{DeviceGlobalProp->Name}; + assert(DeviceGlobalName.find(Prefix) == 0); + bool Inserted = false; + std::tie(std::ignore, Inserted) = + DeviceGlobalIDSet.emplace(DeviceGlobalName); + if (Inserted) { + DeviceGlobalIDVec.emplace_back(DeviceGlobalName); + DeviceGlobalName.remove_prefix(Prefix.length()); + DeviceGlobalNames.emplace_back(DeviceGlobalName); + } + } + } + + // Device globals are usually statically allocated and registered in the + // integration footer, which we don't have in the RTC context. Instead, we + // dynamically allocate storage tied to the executable kernel bundle. + for (DeviceGlobalMapEntry *DeviceGlobalEntry : + PM.getDeviceGlobalEntries(DeviceGlobalIDVec)) { + + size_t AllocSize = DeviceGlobalEntry->MDeviceGlobalTSize; // init value + if (!DeviceGlobalEntry->MIsDeviceImageScopeDecorated) { + // Consider storage for device USM pointer. + AllocSize += sizeof(void *); + } + auto Alloc = std::make_unique(AllocSize); + std::string_view DeviceGlobalName{DeviceGlobalEntry->MUniqueId}; + PM.addOrInitDeviceGlobalEntry(Alloc.get(), DeviceGlobalName.data()); + DeviceGlobalAllocations.push_back(std::move(Alloc)); + + // Drop the RTC prefix from the entry's symbol name. Note that the PM + // still manages this device global under its prefixed name. + assert(DeviceGlobalName.find(Prefix) == 0); + DeviceGlobalName.remove_prefix(Prefix.length()); + DeviceGlobalEntry->MUniqueId = DeviceGlobalName; } return std::make_shared( MContext, MDevices, KernelIDs, std::move(KernelNames), - std::move(MangledKernelNames), Binaries, std::move(Prefix), + std::move(MangledKernelNames), std::move(DeviceGlobalNames), + std::move(DeviceGlobalAllocations), Binaries, std::move(Prefix), MLanguage); } @@ -680,6 +729,8 @@ class kernel_bundle_impl { KernelNames, MLanguage); } + // Utility methods for kernel_compiler functionality +private: std::string adjust_kernel_name(const std::string &Name) { if (MLanguage == syclex::source_language::sycl) { auto It = MMangledKernelNames.find(Name); @@ -694,8 +745,58 @@ class kernel_bundle_impl { MKernelNames.end(); } + std::string mangle_device_global_name(const std::string &Name) { + // TODO: Support device globals declared in namespaces. + return "_Z" + std::to_string(Name.length()) + Name; + } + + DeviceGlobalMapEntry *get_device_global_entry(const std::string &Name) { + if (MKernelNames.empty() || MLanguage != syclex::source_language::sycl) { + throw sycl::exception(make_error_code(errc::invalid), + "Querying device globals by name is only available " + "in kernel_bundles successfully built from " + "kernel_bundle::ext_oneapi_source> " + "with 'sycl' source language."); + } + + if (!ext_oneapi_has_device_global(Name)) { + throw sycl::exception(make_error_code(errc::invalid), + "device global '" + Name + + "' not found in kernel_bundle"); + } + + std::vector Entries = + ProgramManager::getInstance().getDeviceGlobalEntries( + {MPrefix + mangle_device_global_name(Name)}); + assert(Entries.size() == 1); + return Entries.front(); + } + + void unregister_device_globals_from_context() { + if (MDeviceGlobalNames.empty()) + return; + + // Manually trigger the release of resources for all device global map + // entries associated with this runtime-compiled bundle. Normally, this + // would happen in `~context_impl()`, however in the RTC setting, the + // context outlives the DG map entries owned by the program manager. + + std::vector DeviceGlobalIDs; + std::transform(MDeviceGlobalNames.begin(), MDeviceGlobalNames.end(), + std::back_inserter(DeviceGlobalIDs), + [&](const std::string &DGName) { return MPrefix + DGName; }); + auto ContextImpl = getSyclObjImpl(MContext); + for (DeviceGlobalMapEntry *Entry : + ProgramManager::getInstance().getDeviceGlobalEntries( + DeviceGlobalIDs)) { + Entry->removeAssociatedResources(ContextImpl.get()); + ContextImpl->removeAssociatedDeviceGlobal(Entry->MDeviceGlobalPtr); + } + } + +public: bool ext_oneapi_has_kernel(const std::string &Name) { - return is_kernel_name(adjust_kernel_name(Name)); + return !MKernelNames.empty() && is_kernel_name(adjust_kernel_name(Name)); } kernel @@ -768,6 +869,41 @@ class kernel_bundle_impl { return AdjustedName; } + bool ext_oneapi_has_device_global(const std::string &Name) { + return !MDeviceGlobalNames.empty() && + std::find(MDeviceGlobalNames.begin(), MDeviceGlobalNames.end(), + mangle_device_global_name(Name)) != + MDeviceGlobalNames.end(); + } + + void *ext_oneapi_get_device_global_address(const std::string &Name, + const device &Dev) { + DeviceGlobalMapEntry *Entry = get_device_global_entry(Name); + + if (std::find(MDevices.begin(), MDevices.end(), Dev) == MDevices.end()) { + throw sycl::exception(make_error_code(errc::invalid), + "kernel_bundle not built for device"); + } + + if (Entry->MIsDeviceImageScopeDecorated) { + throw sycl::exception(make_error_code(errc::invalid), + "Cannot query USM pointer for device global with " + "'device_image_scope' property"); + } + + // TODO: Add context-only initialization via `urUSMContextMemcpyExp` instead + // of using a throw-away queue. + queue InitQueue{MContext, Dev}; + auto &USMMem = + Entry->getOrAllocateDeviceGlobalUSM(getSyclObjImpl(InitQueue)); + InitQueue.wait_and_throw(); + return USMMem.getPtr(); + } + + size_t ext_oneapi_get_device_global_size(const std::string &Name) { + return get_device_global_entry(Name)->MDeviceGlobalTSize; + } + bool empty() const noexcept { return MDeviceImages.empty(); } backend get_backend() const noexcept { @@ -999,6 +1135,7 @@ class kernel_bundle_impl { ~kernel_bundle_impl() { try { if (MDeviceBinaries) { + unregister_device_globals_from_context(); ProgramManager::getInstance().removeImages(MDeviceBinaries); syclex::detail::SYCL_JIT_Destroy(MDeviceBinaries); } @@ -1039,6 +1176,8 @@ class kernel_bundle_impl { // only kernel_bundles created from source have KernelNames member. std::vector MKernelNames; std::unordered_map MMangledKernelNames; + std::vector MDeviceGlobalNames; + std::vector> MDeviceGlobalAllocations; sycl_device_binaries MDeviceBinaries = nullptr; std::string MPrefix; include_pairs_t MIncludePairs; diff --git a/sycl/source/kernel_bundle.cpp b/sycl/source/kernel_bundle.cpp index 0ffb764bde796..db899bee36913 100644 --- a/sycl/source/kernel_bundle.cpp +++ b/sycl/source/kernel_bundle.cpp @@ -142,6 +142,21 @@ 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())}; } +bool kernel_bundle_plain::ext_oneapi_has_device_global( + detail::string_view name) { + return impl->ext_oneapi_has_device_global(name.data()); +} + +void *kernel_bundle_plain::ext_oneapi_get_device_global_address( + detail::string_view name, const device &dev) { + return impl->ext_oneapi_get_device_global_address(name.data(), dev); +} + +size_t kernel_bundle_plain::ext_oneapi_get_device_global_size( + detail::string_view name) { + return impl->ext_oneapi_get_device_global_size(name.data()); +} + ////////////////////////////////// ///// sycl::detail free functions ////////////////////////////////// diff --git a/sycl/test-e2e/KernelCompiler/sycl_device_globals.cpp b/sycl/test-e2e/KernelCompiler/sycl_device_globals.cpp new file mode 100644 index 0000000000000..932e472ffa0dd --- /dev/null +++ b/sycl/test-e2e/KernelCompiler/sycl_device_globals.cpp @@ -0,0 +1,173 @@ +//==--- sycl_device_globals.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) +// REQUIRES: aspect-usm_device_allocations + +// UNSUPPORTED: accelerator, opencl && gpu +// UNSUPPORTED-INTENDED: while accelerator is AoT only, this cannot run there. +// UNSUPPORTED: opencl && gpu +// UNSUPPORTED-TRACKER: GSD-4287 + +// RUN: %{build} -o %t.out +// RUN: %{run} %t.out +// RUN: %{l0_leak_check} %{run} %t.out + +#include +#include +#include + +auto constexpr DGSource = R"===( +#include + +namespace syclex = sycl::ext::oneapi::experimental; + +syclex::device_global DG; + +extern "C" SYCL_EXTERNAL SYCL_EXT_ONEAPI_FUNCTION_PROPERTY( + (syclex::single_task_kernel)) void ff_dg_adder(int val) { + DG += val; +} + +syclex::device_global DG_DIS; + +extern "C" SYCL_EXTERNAL SYCL_EXT_ONEAPI_FUNCTION_PROPERTY( + (syclex::single_task_kernel)) void ff_swap(int64_t *val) { + int64_t tmp = DG_DIS; + DG_DIS = *val; + *val = tmp; +} + +)==="; + +int test_device_global() { + 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(); + sycl::device d = q.get_device(); + + bool ok = d.ext_oneapi_can_compile(syclex::source_language::sycl); + if (!ok) { + std::cout << "Apparently this device does not support `sycl` source kernel " + "bundle extension: " + << d.get_info() << std::endl; + return -1; + } + + source_kb kbSrc = syclex::create_kernel_bundle_from_source( + ctx, syclex::source_language::sycl, DGSource); + + exe_kb kbExe1 = syclex::build(kbSrc); + auto addK = kbExe1.ext_oneapi_get_kernel("ff_dg_adder"); + + // Check presence of device globals. + assert(kbExe1.ext_oneapi_has_device_global("DG")); + // Querying a non-existing device global shall not crash. + assert(!kbExe1.ext_oneapi_has_device_global("bogus_DG")); + + void *dgAddr = kbExe1.ext_oneapi_get_device_global_address("DG", d); + size_t dgSize = kbExe1.ext_oneapi_get_device_global_size("DG"); + assert(dgSize == 4); + + int32_t val; + auto checkVal = [&](int32_t expected) { + val = -1; + q.memcpy(&val, dgAddr, dgSize).wait(); + std::cout << "val: " << val << " == " << expected << '\n'; + assert(val == expected); + }; + + // Device globals are zero-initialized. + checkVal(0); + + // Set the DG. + val = 123; + q.memcpy(dgAddr, &val, dgSize).wait(); + checkVal(123); + + // Run a kernel using it. + val = -17; + q.submit([&](sycl::handler &CGH) { + CGH.set_arg(0, val); + CGH.single_task(addK); + }); + q.wait(); + checkVal(123 - 17); + + // Test that each bundle has its distinct set of globals. + exe_kb kbExe2 = syclex::build(kbSrc); + dgAddr = kbExe2.ext_oneapi_get_device_global_address("DG", d); + checkVal(0); + + dgAddr = kbExe1.ext_oneapi_get_device_global_address("DG", d); + checkVal(123 - 17); + + // Test global with `device_image_scope`. We currently cannot read/write these + // from the host, but they should work device-only. + auto swapK = kbExe2.ext_oneapi_get_kernel("ff_swap"); + int64_t *valBuf = sycl::malloc_shared(1, q); + *valBuf = -1; + auto doSwap = [&]() { + q.submit([&](sycl::handler &CGH) { + CGH.set_arg(0, valBuf); + CGH.single_task(swapK); + }); + q.wait(); + }; + + doSwap(); + assert(*valBuf == 0); + doSwap(); + assert(*valBuf == -1); + + sycl::free(valBuf, q); + + return 0; +} + +int 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(); + sycl::device d = q.get_device(); + + bool ok = d.ext_oneapi_can_compile(syclex::source_language::sycl); + if (!ok) { + return 0; + } + + source_kb kbSrc = syclex::create_kernel_bundle_from_source( + ctx, syclex::source_language::sycl, DGSource); + exe_kb kbExe = syclex::build(kbSrc); + + try { + kbExe.ext_oneapi_get_device_global_address("DG_DIS", d); + assert(false && "we should not be here"); + } catch (sycl::exception &e) { + assert(e.code() == sycl::errc::invalid); + assert(std::string(e.what()).find( + "Cannot query USM pointer for device global with " + "'device_image_scope' property") != std::string::npos); + } + return 0; +} + +int main() { +#ifdef SYCL_EXT_ONEAPI_KERNEL_COMPILER + return test_device_global() || test_error(); +#else + static_assert(false, "Kernel Compiler feature test macro undefined"); +#endif + return 0; +} diff --git a/sycl/test/abi/sycl_symbols_linux.dump b/sycl/test/abi/sycl_symbols_linux.dump index cbe08c9c12d4f..564d64026ecc1 100644 --- a/sycl/test/abi/sycl_symbols_linux.dump +++ b/sycl/test/abi/sycl_symbols_linux.dump @@ -3278,8 +3278,11 @@ _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_plain28ext_oneapi_has_device_globalENS1_11string_viewE _ZN4sycl3_V16detail19kernel_bundle_plain30ext_oneapi_get_raw_kernel_nameENS1_11string_viewE _ZN4sycl3_V16detail19kernel_bundle_plain32set_specialization_constant_implEPKcPvm +_ZN4sycl3_V16detail19kernel_bundle_plain33ext_oneapi_get_device_global_sizeENS1_11string_viewE +_ZN4sycl3_V16detail19kernel_bundle_plain36ext_oneapi_get_device_global_addressENS1_11string_viewERKNS0_6deviceE _ZN4sycl3_V16detail20associateWithHandlerERNS0_7handlerEPNS1_16AccessorBaseHostENS0_6access6targetE _ZN4sycl3_V16detail20associateWithHandlerERNS0_7handlerEPNS1_28SampledImageAccessorBaseHostENS0_12image_targetE _ZN4sycl3_V16detail20associateWithHandlerERNS0_7handlerEPNS1_30UnsampledImageAccessorBaseHostENS0_12image_targetE diff --git a/sycl/test/abi/sycl_symbols_windows.dump b/sycl/test/abi/sycl_symbols_windows.dump index 06907b875aac5..c55a840b78ada 100644 --- a/sycl/test/abi/sycl_symbols_windows.dump +++ b/sycl/test/abi/sycl_symbols_windows.dump @@ -3884,6 +3884,10 @@ ?ext_oneapi_get_composite_devices@platform@_V1@sycl@@QEBA?AV?$vector@Vdevice@_V1@sycl@@V?$allocator@Vdevice@_V1@sycl@@@std@@@std@@XZ ?ext_oneapi_get_default_context@platform@_V1@sycl@@QEBA?AVcontext@23@XZ ?khr_get_default_context@platform@_V1@sycl@@QEBA?AVcontext@23@XZ +?ext_oneapi_get_device_global_address@kernel_bundle_plain@detail@_V1@sycl@@AEAAPEAXVstring_view@234@AEBVdevice@34@@Z +?ext_oneapi_get_device_global_address@kernel_bundle_plain@detail@_V1@sycl@@QEAAPEAXAEBV?$basic_string@DU?$char_traits@D@std@@V?$allocator@D@2@@std@@AEBVdevice@34@@Z +?ext_oneapi_get_device_global_size@kernel_bundle_plain@detail@_V1@sycl@@AEAA_KVstring_view@234@@Z +?ext_oneapi_get_device_global_size@kernel_bundle_plain@detail@_V1@sycl@@QEAA_KAEBV?$basic_string@DU?$char_traits@D@std@@V?$allocator@D@2@@std@@@Z ?ext_oneapi_get_graph@queue@_V1@sycl@@QEBA?AV?$command_graph@$0A@@experimental@oneapi@ext@23@XZ ?ext_oneapi_get_kernel@kernel_bundle_plain@detail@_V1@sycl@@AEAA?AVkernel@34@Vstring_view@234@@Z ?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 @@ -3896,6 +3900,8 @@ ?ext_oneapi_graph@queue@_V1@sycl@@QEAA?AVevent@23@V?$command_graph@$00@experimental@oneapi@ext@23@AEBUcode_location@detail@23@@Z ?ext_oneapi_graph@queue@_V1@sycl@@QEAA?AVevent@23@V?$command_graph@$00@experimental@oneapi@ext@23@AEBV?$vector@Vevent@_V1@sycl@@V?$allocator@Vevent@_V1@sycl@@@std@@@std@@AEBUcode_location@detail@23@@Z ?ext_oneapi_graph@queue@_V1@sycl@@QEAA?AVevent@23@V?$command_graph@$00@experimental@oneapi@ext@23@V423@AEBUcode_location@detail@23@@Z +?ext_oneapi_has_device_global@kernel_bundle_plain@detail@_V1@sycl@@AEAA_NVstring_view@234@@Z +?ext_oneapi_has_device_global@kernel_bundle_plain@detail@_V1@sycl@@QEAA_NAEBV?$basic_string@DU?$char_traits@D@std@@V?$allocator@D@2@@std@@@Z ?ext_oneapi_has_kernel@kernel_bundle_plain@detail@_V1@sycl@@AEAA_NVstring_view@234@@Z ?ext_oneapi_has_kernel@kernel_bundle_plain@detail@_V1@sycl@@QEAA_NAEBV?$basic_string@DU?$char_traits@D@std@@V?$allocator@D@2@@std@@@Z ?ext_oneapi_memcpy2d_impl@handler@_V1@sycl@@AEAAXPEAX_KPEBX111@Z diff --git a/sycl/test/e2e_test_requirements/no_sycl_hpp_in_e2e_tests.cpp b/sycl/test/e2e_test_requirements/no_sycl_hpp_in_e2e_tests.cpp index cecce9368f21b..e51d5f86e1777 100644 --- a/sycl/test/e2e_test_requirements/no_sycl_hpp_in_e2e_tests.cpp +++ b/sycl/test/e2e_test_requirements/no_sycl_hpp_in_e2e_tests.cpp @@ -6,7 +6,7 @@ // CHECK-DAG: README.md // CHECK-DAG: lit.cfg.py // -// CHECK-NUM-MATCHES: 11 +// CHECK-NUM-MATCHES: 12 // // This test verifies that `` isn't used in E2E tests. Instead, // fine-grained includes should used, see