diff --git a/sycl/include/CL/sycl/accessor.hpp b/sycl/include/CL/sycl/accessor.hpp index 423785dc928be..26a1624c56d46 100644 --- a/sycl/include/CL/sycl/accessor.hpp +++ b/sycl/include/CL/sycl/accessor.hpp @@ -1195,7 +1195,9 @@ class __SYCL_SPECIAL_CLASS accessor : buffer &BufferRef, TagT, const property_list &PropertyList = {}, const detail::code_location CodeLoc = detail::code_location::current()) - : accessor(BufferRef, PropertyList, CodeLoc) {} + : accessor(BufferRef, PropertyList, CodeLoc) { + adjustAccPropsInBuf(detail::getSyclObjImpl(BufferRef).get()); + } template &PropertyList = {}, const detail::code_location CodeLoc = detail::code_location::current()) - : accessor(BufferRef, PropertyList, CodeLoc) {} + : accessor(BufferRef, PropertyList, CodeLoc) { + adjustAccPropsInBuf(detail::getSyclObjImpl(BufferRef).get()); + } #endif template &BufferRef, handler &CommandGroupHandler, TagT, const property_list &PropertyList = {}, const detail::code_location CodeLoc = detail::code_location::current()) - : accessor(BufferRef, CommandGroupHandler, PropertyList, CodeLoc) {} + : accessor(BufferRef, CommandGroupHandler, PropertyList, CodeLoc) { + adjustAccPropsInBuf(detail::getSyclObjImpl(BufferRef).get()); + } template &PropertyList = {}, const detail::code_location CodeLoc = detail::code_location::current()) - : accessor(BufferRef, CommandGroupHandler, PropertyList, CodeLoc) {} + : accessor(BufferRef, CommandGroupHandler, PropertyList, CodeLoc) { + adjustAccPropsInBuf(detail::getSyclObjImpl(BufferRef).get()); + } #endif @@ -1341,7 +1349,9 @@ class __SYCL_SPECIAL_CLASS accessor : buffer &BufferRef, range AccessRange, TagT, const property_list &PropertyList = {}, const detail::code_location CodeLoc = detail::code_location::current()) - : accessor(BufferRef, AccessRange, {}, PropertyList, CodeLoc) {} + : accessor(BufferRef, AccessRange, {}, PropertyList, CodeLoc) { + adjustAccPropsInBuf(detail::getSyclObjImpl(BufferRef).get()); + } template &PropertyList = {}, const detail::code_location CodeLoc = detail::code_location::current()) - : accessor(BufferRef, AccessRange, {}, PropertyList, CodeLoc) {} + : accessor(BufferRef, AccessRange, {}, PropertyList, CodeLoc) { + adjustAccPropsInBuf(detail::getSyclObjImpl(BufferRef).get()); + } #endif template &BufferRef, range AccessRange, id AccessOffset, TagT, const property_list &PropertyList = {}, const detail::code_location CodeLoc = detail::code_location::current()) - : accessor(BufferRef, AccessRange, AccessOffset, PropertyList, CodeLoc) {} + : accessor(BufferRef, AccessRange, AccessOffset, PropertyList, CodeLoc) { + adjustAccPropsInBuf(detail::getSyclObjImpl(BufferRef).get()); + } template &PropertyList = {}, const detail::code_location CodeLoc = detail::code_location::current()) - : accessor(BufferRef, AccessRange, AccessOffset, PropertyList, CodeLoc) {} + : accessor(BufferRef, AccessRange, AccessOffset, PropertyList, CodeLoc) { + adjustAccPropsInBuf(detail::getSyclObjImpl(BufferRef).get()); + } #endif template @@ -1802,6 +1826,32 @@ class __SYCL_SPECIAL_CLASS accessor : PI_INVALID_VALUE); } } + +#if __cplusplus >= 201703L + template + void adjustAccPropsInBuf(detail::SYCLMemObjI *SYCLMemObject) { + if constexpr (PropertyListT::template has_property< + sycl::ext::intel::property::buffer_location>()) { + auto location = (PropertyListT::template get_property< + sycl::ext::intel::property::buffer_location>()) + .get_location(); + property_list PropList{ + sycl::property::buffer::detail::buffer_location(location)}; + detail::SYCLMemObjT *SYCLMemObjectT = + dynamic_cast(SYCLMemObject); + SYCLMemObjectT->addOrReplaceAccessorProperties(PropList); + } else { + deleteAccPropsFromBuf(SYCLMemObject); + } + } + + void deleteAccPropsFromBuf(detail::SYCLMemObjI *SYCLMemObject) { + detail::SYCLMemObjT *SYCLMemObjectT = + dynamic_cast(SYCLMemObject); + SYCLMemObjectT->deleteAccessorProperty( + sycl::detail::PropWithDataKind::AccPropBufferLocation); + } +#endif }; #if __cplusplus >= 201703L diff --git a/sycl/include/CL/sycl/detail/pi.h b/sycl/include/CL/sycl/detail/pi.h index 624078b260687..de86580d18abc 100644 --- a/sycl/include/CL/sycl/detail/pi.h +++ b/sycl/include/CL/sycl/detail/pi.h @@ -591,6 +591,8 @@ constexpr pi_map_flags PI_MAP_WRITE_INVALIDATE_REGION = // make the translation to OpenCL transparent. using pi_mem_properties = pi_bitfield; constexpr pi_mem_properties PI_MEM_PROPERTIES_CHANNEL = CL_MEM_CHANNEL_INTEL; +constexpr pi_mem_properties PI_MEM_PROPERTIES_ALLOC_BUFFER_LOCATION = + CL_MEM_ALLOC_BUFFER_LOCATION_INTEL; // NOTE: this is made 64-bit to match the size of cl_mem_properties_intel to // make the translation to OpenCL transparent. diff --git a/sycl/include/CL/sycl/detail/property_helper.hpp b/sycl/include/CL/sycl/detail/property_helper.hpp index ecd942f310201..291327c13caab 100644 --- a/sycl/include/CL/sycl/detail/property_helper.hpp +++ b/sycl/include/CL/sycl/detail/property_helper.hpp @@ -48,7 +48,7 @@ enum PropWithDataKind { ImageContextBound = 3, BufferMemChannel = 4, AccPropBufferLocation = 5, - PropWithDataKindSize = 6 + PropWithDataKindSize = 6, }; // Base class for dataless properties, needed to check that the type of an diff --git a/sycl/include/CL/sycl/detail/property_list_base.hpp b/sycl/include/CL/sycl/detail/property_list_base.hpp index 0bdea3e85b7ae..eed8cf9ad4252 100644 --- a/sycl/include/CL/sycl/detail/property_list_base.hpp +++ b/sycl/include/CL/sycl/detail/property_list_base.hpp @@ -103,6 +103,30 @@ class PropertyListBase { PI_INVALID_VALUE); } + void add_or_replace_accessor_properties_helper( + const std::vector> &PropsWithData) { + for (auto &Prop : PropsWithData) { + if (Prop->isSame(sycl::detail::PropWithDataKind::AccPropBufferLocation)) { + delete_accessor_property_helper( + sycl::detail::PropWithDataKind::AccPropBufferLocation); + MPropsWithData.push_back(Prop); + break; + } + } + } + + void delete_accessor_property_helper(const PropWithDataKind &Kind) { + auto It = MPropsWithData.begin(); + for (; It != MPropsWithData.end(); ++It) { + if ((*It)->isSame(Kind)) + break; + } + if (It != MPropsWithData.end()) { + std::iter_swap(It, MPropsWithData.end() - 1); + MPropsWithData.pop_back(); + } + } + // Stores enabled/disabled for simple properties std::bitset MDataLessProps; // Stores shared_ptrs to complex properties diff --git a/sycl/include/CL/sycl/detail/sycl_mem_obj_t.hpp b/sycl/include/CL/sycl/detail/sycl_mem_obj_t.hpp index 09639d011f923..bd26c70a1d972 100644 --- a/sycl/include/CL/sycl/detail/sycl_mem_obj_t.hpp +++ b/sycl/include/CL/sycl/detail/sycl_mem_obj_t.hpp @@ -111,6 +111,15 @@ class __SYCL_EXPORT SYCLMemObjT : public SYCLMemObjI { return MProps.get_property(); } + __SYCL_DLL_LOCAL void + addOrReplaceAccessorProperties(const property_list &PropertyList) { + MProps.add_or_replace_accessor_properties(PropertyList); + } + + __SYCL_DLL_LOCAL void deleteAccessorProperty(const PropWithDataKind &Kind) { + MProps.delete_accessor_property(Kind); + } + template __SYCL_DLL_LOCAL AllocatorT get_allocator() const { return MAllocator->getAllocator(); diff --git a/sycl/include/CL/sycl/properties/accessor_properties.hpp b/sycl/include/CL/sycl/properties/accessor_properties.hpp index 4b608ff9c2fce..a0a97649c7cba 100644 --- a/sycl/include/CL/sycl/properties/accessor_properties.hpp +++ b/sycl/include/CL/sycl/properties/accessor_properties.hpp @@ -58,6 +58,7 @@ struct buffer_location { constexpr bool operator!=(const buffer_location::instance &) const { return A != B; } + int get_location() { return A; } }; }; } // namespace property diff --git a/sycl/include/CL/sycl/properties/buffer_properties.hpp b/sycl/include/CL/sycl/properties/buffer_properties.hpp index fc36a17b77ca5..15203cda547c3 100644 --- a/sycl/include/CL/sycl/properties/buffer_properties.hpp +++ b/sycl/include/CL/sycl/properties/buffer_properties.hpp @@ -51,6 +51,18 @@ class mem_channel : public detail::PropertyWithData< uint32_t MChannel; }; +namespace detail { +class buffer_location + : public sycl::detail::PropertyWithData< + sycl::detail::PropWithDataKind::AccPropBufferLocation> { +public: + buffer_location(uint64_t Location) : MLocation(Location) {} + uint64_t get_buffer_location() const { return MLocation; } + +private: + uint64_t MLocation; +}; +} // namespace detail } // namespace buffer } // namespace property @@ -75,6 +87,9 @@ template <> struct is_property : std::true_type {}; template <> struct is_property : std::true_type {}; template <> +struct is_property : std::true_type { +}; +template <> struct is_property : std::true_type {}; template <> struct is_property : std::true_type {}; @@ -91,6 +106,10 @@ struct is_property_of> : std::true_type {}; template +struct is_property_of> + : std::true_type {}; +template struct is_property_of> : std::true_type {}; diff --git a/sycl/include/CL/sycl/property_list.hpp b/sycl/include/CL/sycl/property_list.hpp index fa523a6676e2a..8c9a8183a6c1f 100644 --- a/sycl/include/CL/sycl/property_list.hpp +++ b/sycl/include/CL/sycl/property_list.hpp @@ -51,6 +51,13 @@ class property_list : protected detail::PropertyListBase { return has_property_helper(); } + void add_or_replace_accessor_properties(const property_list &PropertyList) { + add_or_replace_accessor_properties_helper(PropertyList.MPropsWithData); + } + void delete_accessor_property(const sycl::detail::PropWithDataKind &Kind) { + delete_accessor_property_helper(Kind); + } + template operator ext::oneapi::accessor_property_list(); private: diff --git a/sycl/source/detail/context_impl.cpp b/sycl/source/detail/context_impl.cpp index 18390606df128..c10188b83fc15 100644 --- a/sycl/source/detail/context_impl.cpp +++ b/sycl/source/detail/context_impl.cpp @@ -28,7 +28,8 @@ namespace detail { context_impl::context_impl(const device &Device, async_handler AsyncHandler, const property_list &PropList) : MAsyncHandler(AsyncHandler), MDevices(1, Device), MContext(nullptr), - MPlatform(), MPropList(PropList), MHostContext(Device.is_host()) { + MPlatform(), MPropList(PropList), MHostContext(Device.is_host()), + MSupportBufferLocationByDevices(NotChecked) { MKernelProgramCache.setContextPtr(this); } @@ -36,7 +37,8 @@ context_impl::context_impl(const std::vector Devices, async_handler AsyncHandler, const property_list &PropList) : MAsyncHandler(AsyncHandler), MDevices(Devices), MContext(nullptr), - MPlatform(), MPropList(PropList), MHostContext(false) { + MPlatform(), MPropList(PropList), MHostContext(false), + MSupportBufferLocationByDevices(NotChecked) { MPlatform = detail::getSyclObjImpl(MDevices[0].get_platform()); std::vector DeviceIds; for (const auto &D : MDevices) { @@ -66,7 +68,7 @@ context_impl::context_impl(const std::vector Devices, context_impl::context_impl(RT::PiContext PiContext, async_handler AsyncHandler, const plugin &Plugin) : MAsyncHandler(AsyncHandler), MDevices(), MContext(PiContext), MPlatform(), - MHostContext(false) { + MHostContext(false), MSupportBufferLocationByDevices(NotChecked) { std::vector DeviceIds; size_t DevicesNum = 0; @@ -206,6 +208,20 @@ pi_native_handle context_impl::getNative() const { return Handle; } +bool context_impl::isBufferLocationSupported() const { + if (MSupportBufferLocationByDevices != NotChecked) + return MSupportBufferLocationByDevices == Supported ? true : false; + // Check that devices within context have support of buffer location + MSupportBufferLocationByDevices = Supported; + for (auto &Device : MDevices) { + if (!Device.has_extension("cl_intel_mem_alloc_buffer_location")) { + MSupportBufferLocationByDevices = NotSupported; + break; + } + } + return MSupportBufferLocationByDevices == Supported ? true : false; +} + } // namespace detail } // namespace sycl } // __SYCL_INLINE_NAMESPACE(cl) diff --git a/sycl/source/detail/context_impl.hpp b/sycl/source/detail/context_impl.hpp index d1a52181e9d33..9049ab1a12760 100644 --- a/sycl/source/detail/context_impl.hpp +++ b/sycl/source/detail/context_impl.hpp @@ -167,6 +167,11 @@ class context_impl { /// \return a native handle. pi_native_handle getNative() const; + // Returns true if buffer_location property is supported by devices + bool isBufferLocationSupported() const; + + enum PropertySupport { NotSupported = 0, Supported = 1, NotChecked = 2 }; + private: async_handler MAsyncHandler; std::vector MDevices; @@ -177,6 +182,7 @@ class context_impl { std::map, RT::PiProgram> MCachedLibPrograms; mutable KernelProgramCache MKernelProgramCache; + mutable PropertySupport MSupportBufferLocationByDevices; }; } // namespace detail diff --git a/sycl/source/detail/memory_manager.cpp b/sycl/source/detail/memory_manager.cpp index 1cd77083e0c16..2a966439776ca 100644 --- a/sycl/source/detail/memory_manager.cpp +++ b/sycl/source/detail/memory_manager.cpp @@ -362,6 +362,18 @@ MemoryManager::allocateBufferObject(ContextImplPtr TargetContext, void *UserPtr, RT::PiMem NewMem = nullptr; const detail::plugin &Plugin = TargetContext->getPlugin(); + + if (PropsList.has_property()) + if (TargetContext->isBufferLocationSupported()) { + auto location = + PropsList.get_property() + .get_buffer_location(); + pi_mem_properties props[3] = {PI_MEM_PROPERTIES_ALLOC_BUFFER_LOCATION, + location, 0}; + memBufferCreateHelper(Plugin, TargetContext->getHandleRef(), + CreationFlags, Size, UserPtr, &NewMem, props); + return NewMem; + } memBufferCreateHelper(Plugin, TargetContext->getHandleRef(), CreationFlags, Size, UserPtr, &NewMem, nullptr); return NewMem; diff --git a/sycl/test/abi/sycl_symbols_windows.dump b/sycl/test/abi/sycl_symbols_windows.dump index 4c6d0e99c0c00..146c57b18a6c3 100644 --- a/sycl/test/abi/sycl_symbols_windows.dump +++ b/sycl/test/abi/sycl_symbols_windows.dump @@ -1088,6 +1088,7 @@ ?acospi@__host_std@cl@@YAMM@Z ?acospi@__host_std@cl@@YANN@Z ?addHostAccessorAndWait@detail@sycl@cl@@YAXPEAVAccessorImplHost@123@@Z +?addOrReplaceAccessorProperties@SYCLMemObjT@detail@sycl@cl@@QEAAXAEBVproperty_list@34@@Z ?addReduction@handler@sycl@cl@@AEAAXAEBV?$shared_ptr@$$CBX@std@@@Z ?addStream@handler@sycl@cl@@AEAAXAEBV?$shared_ptr@Vstream_impl@detail@sycl@cl@@@std@@@Z ?advise_usm@MemoryManager@detail@sycl@cl@@SAXPEBXV?$shared_ptr@Vqueue_impl@detail@sycl@cl@@@std@@_KW4_pi_mem_advice@@V?$vector@PEAU_pi_event@@V?$allocator@PEAU_pi_event@@@std@@@6@AEAPEAU_pi_event@@@Z @@ -1686,6 +1687,7 @@ ?degrees@__host_std@cl@@YA?AVhalf@half_impl@detail@sycl@2@V34562@@Z ?degrees@__host_std@cl@@YAMM@Z ?degrees@__host_std@cl@@YANN@Z +?deleteAccessorProperty@SYCLMemObjT@detail@sycl@cl@@QEAAXAEBW4PropWithDataKind@234@@Z ?depends_on@handler@sycl@cl@@QEAAXAEBV?$vector@Vevent@sycl@cl@@V?$allocator@Vevent@sycl@cl@@@std@@@std@@@Z ?depends_on@handler@sycl@cl@@QEAAXVevent@23@@Z ?destructorNotification@buffer_impl@detail@sycl@cl@@QEAAXPEAX@Z diff --git a/sycl/unittests/CMakeLists.txt b/sycl/unittests/CMakeLists.txt index 59e7979dd03e7..83eb6cccc44b7 100644 --- a/sycl/unittests/CMakeLists.txt +++ b/sycl/unittests/CMakeLists.txt @@ -34,3 +34,4 @@ add_subdirectory(assert) add_subdirectory(Extensions) add_subdirectory(windows) add_subdirectory(event) +add_subdirectory(buffer) diff --git a/sycl/unittests/buffer/BufferLocation.cpp b/sycl/unittests/buffer/BufferLocation.cpp new file mode 100644 index 0000000000000..5d2313ee5dd15 --- /dev/null +++ b/sycl/unittests/buffer/BufferLocation.cpp @@ -0,0 +1,216 @@ +//==-------- buffer_location.cpp --- check buffer_location property --------==// +// +// 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 +// +//===----------------------------------------------------------------------===// +#define SYCL2020_DISABLE_DEPRECATION_WARNINGS + +#include +#include +#include + +#include +#include + +#include + +const uint64_t DEFAULT_VALUE = 7777; +static uint64_t PassedLocation = DEFAULT_VALUE; + +pi_result redefinedMemBufferCreate(pi_context, pi_mem_flags, size_t size, + void *, pi_mem *, + const pi_mem_properties *properties) { + PassedLocation = DEFAULT_VALUE; + if (!properties) + return PI_SUCCESS; + + // properties must ended by 0 + size_t I = 0; + while (true) { + if (properties[I] != 0) { + if (properties[I] != PI_MEM_PROPERTIES_ALLOC_BUFFER_LOCATION) { + I += 2; + } else { + PassedLocation = properties[I + 1]; + break; + } + } + } + + return PI_SUCCESS; +} + +static pi_result redefinedDeviceGetInfo(pi_device device, + pi_device_info param_name, + size_t param_value_size, + void *param_value, + size_t *param_value_size_ret) { + if (param_name == PI_DEVICE_INFO_TYPE) { + auto *Result = reinterpret_cast<_pi_device_type *>(param_value); + *Result = PI_DEVICE_TYPE_ACC; + } + if (param_name == PI_DEVICE_INFO_COMPILER_AVAILABLE) { + auto *Result = reinterpret_cast(param_value); + *Result = true; + } + if (param_name == PI_DEVICE_INFO_EXTENSIONS) { + const std::string name = "cl_intel_mem_alloc_buffer_location"; + if (!param_value) { + *param_value_size_ret = name.size(); + } else { + char *dst = static_cast(param_value); + strcpy(dst, name.data()); + } + } + return PI_SUCCESS; +} + +class BufferTest : public ::testing::Test { +public: + BufferTest() : Plt{sycl::default_selector()} {} + +protected: + void SetUp() override { + if (Plt.is_host() || Plt.get_backend() != sycl::backend::opencl) { + std::cout << "This test is only supported on OpenCL backend\n"; + std::cout << "Current platform is " + << Plt.get_info(); + return; + } + + Mock = std::make_unique(Plt); + + setupDefaultMockAPIs(*Mock); + Mock->redefine( + redefinedMemBufferCreate); + Mock->redefine( + redefinedDeviceGetInfo); + } + +protected: + std::unique_ptr Mock; + sycl::platform Plt; +}; + +// Test that buffer_location was passed correctly +TEST_F(BufferTest, BufferLocationOnly) { + if (Plt.is_host() || Plt.get_backend() != sycl::backend::opencl) { + return; + } + + sycl::context Context{Plt}; + sycl::queue Queue{Context, sycl::accelerator_selector{}}; + + cl::sycl::buffer Buf(3); + Queue + .submit([&](cl::sycl::handler &cgh) { + sycl::ext::oneapi::accessor_property_list< + cl::sycl::ext::intel::property::buffer_location::instance<2>> + PL{sycl::ext::intel::buffer_location<2>}; + sycl::accessor< + int, 1, cl::sycl::access::mode::read_write, + cl::sycl::access::target::global_buffer, + cl::sycl::access::placeholder::false_t, + cl::sycl::ext::oneapi::accessor_property_list< + cl::sycl::ext::intel::property::buffer_location::instance<2>>> + Acc{Buf, cgh, sycl::read_write, PL}; + cgh.single_task([=]() { Acc[0] = 4; }); + }) + .wait(); + EXPECT_EQ(PassedLocation, (uint64_t)2); +} + +// Test that buffer_location was passed correcty if there is one more accessor +// property and buffer_location is correctly chaned by creating new accessors +TEST_F(BufferTest, BufferLocationWithAnotherProp) { + if (Plt.is_host() || Plt.get_backend() != sycl::backend::opencl) { + return; + } + + sycl::context Context{Plt}; + sycl::queue Queue{Context, sycl::accelerator_selector{}}; + + cl::sycl::buffer Buf(3); + Queue + .submit([&](cl::sycl::handler &cgh) { + sycl::ext::oneapi::accessor_property_list< + cl::sycl::ext::oneapi::property::no_alias::instance, + cl::sycl::ext::intel::property::buffer_location::instance<5>> + PL{sycl::ext::oneapi::no_alias, + sycl::ext::intel::buffer_location<5>}; + sycl::accessor< + int, 1, cl::sycl::access::mode::write, + cl::sycl::access::target::global_buffer, + cl::sycl::access::placeholder::false_t, + cl::sycl::ext::oneapi::accessor_property_list< + cl::sycl::ext::oneapi::property::no_alias::instance, + cl::sycl::ext::intel::property::buffer_location::instance<5>>> + Acc{Buf, cgh, sycl::write_only, PL}; + + cgh.single_task([=]() { Acc[0] = 4; }); + }) + .wait(); + EXPECT_EQ(PassedLocation, (uint64_t)5); + + // Check that if new accessor created, buffer_location is changed + Queue + .submit([&](cl::sycl::handler &cgh) { + sycl::ext::oneapi::accessor_property_list< + cl::sycl::ext::intel::property::buffer_location::instance<3>> + PL{sycl::ext::intel::buffer_location<3>}; + sycl::accessor< + int, 1, cl::sycl::access::mode::write, + cl::sycl::access::target::global_buffer, + cl::sycl::access::placeholder::false_t, + cl::sycl::ext::oneapi::accessor_property_list< + cl::sycl::ext::intel::property::buffer_location::instance<3>>> + Acc{Buf, cgh, sycl::write_only, PL}; + }) + .wait(); + std::shared_ptr BufImpl = + sycl::detail::getSyclObjImpl(Buf); + EXPECT_EQ( + BufImpl->get_property() + .get_buffer_location(), + (uint64_t)3); + + // Check that if new accessor created, buffer_location is deleted from buffer + Queue + .submit([&](cl::sycl::handler &cgh) { + sycl::accessor> + Acc{Buf, cgh, sycl::write_only}; + }) + .wait(); + + EXPECT_EQ( + BufImpl->has_property(), + 0); +} + +// Test that there is no buffer_location property +TEST_F(BufferTest, WOBufferLocation) { + if (Plt.is_host() || Plt.get_backend() != sycl::backend::opencl) { + return; + } + + sycl::context Context{Plt}; + sycl::queue Queue{Context, sycl::accelerator_selector{}}; + + cl::sycl::buffer Buf(3); + Queue + .submit([&](cl::sycl::handler &cgh) { + sycl::accessor> + Acc{Buf, cgh, sycl::read_write}; + cgh.single_task([=]() { Acc[0] = 4; }); + }) + .wait(); + EXPECT_EQ(PassedLocation, DEFAULT_VALUE); +} diff --git a/sycl/unittests/buffer/CMakeLists.txt b/sycl/unittests/buffer/CMakeLists.txt new file mode 100644 index 0000000000000..d8123f31f0982 --- /dev/null +++ b/sycl/unittests/buffer/CMakeLists.txt @@ -0,0 +1,3 @@ +add_sycl_unittest(BufferTests OBJECT + BufferLocation.cpp +)