From 464174d096a26f3eae55246d48a41f7fd4ec342f Mon Sep 17 00:00:00 2001 From: Lorenc Bushi Date: Tue, 14 Jan 2025 13:23:50 -0800 Subject: [PATCH 01/53] Implement backend content extension --- sycl/include/sycl/kernel_bundle.hpp | 27 +++++++++++++++++++++++++++ sycl/source/kernel_bundle.cpp | 14 ++++++++++++++ 2 files changed, 41 insertions(+) diff --git a/sycl/include/sycl/kernel_bundle.hpp b/sycl/include/sycl/kernel_bundle.hpp index a61019efdbf5d..12dda72c1b45c 100644 --- a/sycl/include/sycl/kernel_bundle.hpp +++ b/sycl/include/sycl/kernel_bundle.hpp @@ -37,6 +37,7 @@ #include // for move #include // for hash #include // for vector +#include // for span namespace sycl { inline namespace _V1 { @@ -117,6 +118,12 @@ class __SYCL_EXPORT device_image_plain { protected: detail::DeviceImageImplPtr impl; + backend get_backend() const; + + const std::byte *get_BinaryStart() const; + + const std::byte *get_BinaryEnd() const; + template friend const decltype(Obj::impl) & detail::getSyclObjImpl(const Obj &SyclObject); @@ -145,6 +152,26 @@ class device_image : public detail::device_image_plain, return device_image_plain::has_kernel(KernelID, Dev); } + backend ext_oneapi_get_backend() const noexcept { + return device_image_plain::get_backend(); + } + + template > + std::vector ext_oneapi_get_backend_content() const { + return std::vector(device_image_plain::get_BinaryStart(), + device_image_plain::get_BinaryEnd()); + } + +#ifdef __cpp_lib_span + template > + std::span ext_oneapi_get_content_backend_view() const { + return std::span(device_image_plain::get_BinaryStart(), + device_image_plain::get_BinaryEnd()); + } +#endif + private: device_image(detail::DeviceImageImplPtr Impl) : device_image_plain(std::move(Impl)) {} diff --git a/sycl/source/kernel_bundle.cpp b/sycl/source/kernel_bundle.cpp index e19c2b9df2a75..bfc307c6bb17f 100644 --- a/sycl/source/kernel_bundle.cpp +++ b/sycl/source/kernel_bundle.cpp @@ -45,6 +45,20 @@ ur_native_handle_t device_image_plain::getNative() const { return impl->getNative(); } +backend device_image_plain::get_backend() const { + return impl->get_context().get_backend(); +} + +const std::byte *device_image_plain::get_BinaryStart() const { + return reinterpret_cast( + impl->get_bin_image_ref()->getRawData().BinaryStart); +} + +const std::byte *device_image_plain::get_BinaryEnd() const { + return reinterpret_cast( + impl->get_bin_image_ref()->getRawData().BinaryEnd); +} + //////////////////////////// ///// kernel_bundle_plain /////////////////////////// From 776a3ed3d5726b704f2d214732aacad99d3d2ac0 Mon Sep 17 00:00:00 2001 From: "Bushi, Lorenc" Date: Wed, 15 Jan 2025 13:06:23 -0800 Subject: [PATCH 02/53] Add tests for extension --- ...eapi_device_image_backend_content.asciidoc | 0 .../L0_interop_test.cpp | 87 +++++++++++++++++++ 2 files changed, 87 insertions(+) rename sycl/doc/extensions/{proposed => experimental}/sycl_ext_oneapi_device_image_backend_content.asciidoc (100%) create mode 100644 sycl/test-e2e/DeviceImageBackendContent/L0_interop_test.cpp diff --git a/sycl/doc/extensions/proposed/sycl_ext_oneapi_device_image_backend_content.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_oneapi_device_image_backend_content.asciidoc similarity index 100% rename from sycl/doc/extensions/proposed/sycl_ext_oneapi_device_image_backend_content.asciidoc rename to sycl/doc/extensions/experimental/sycl_ext_oneapi_device_image_backend_content.asciidoc diff --git a/sycl/test-e2e/DeviceImageBackendContent/L0_interop_test.cpp b/sycl/test-e2e/DeviceImageBackendContent/L0_interop_test.cpp new file mode 100644 index 0000000000000..aa77862f08429 --- /dev/null +++ b/sycl/test-e2e/DeviceImageBackendContent/L0_interop_test.cpp @@ -0,0 +1,87 @@ +// REQUIRES: level_zero, level_zero_dev_kit, aspect-usm_shared_allocations +// RUN: %{build} %level_zero_options -fno-sycl-dead-args-optimization -o %t.out +// RUN: %{run} %t.out +// +#include +#include +#include +#include +#include +#include + +namespace syclext = sycl::ext::oneapi; +namespace syclexp = sycl::ext::oneapi::experimental; + +extern"C" SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<1>)) +void iota(int *ptr) { + size_t id = syclext::this_work_item::get_nd_item<1>().get_global_linear_id(); + ptr[id] = 42; +} + +int main() { + sycl::device d([](const sycl::device &d) { + return d.get_backend() == sycl::backend::ext_oneapi_level_zero; + }); + sycl::queue q{d}; + sycl::context ctxt = q.get_context(); + +#ifndef __SYCL_DEVICE_ONLY__ + // First, run the kernel using the SYCL API. + auto Bundle = sycl::get_kernel_bundle(ctxt); + sycl::kernel_id iota_id = syclexp::get_kernel_id(); + sycl::kernel k_iota = Bundle.get_kernel(iota_id); + int *ptr = sycl::malloc_shared(1, q); + q.submit([&](sycl::handler &cgh) { + cgh.set_args(ptr); + cgh.parallel_for(sycl::nd_range{{1}, {1}}, k_iota); + }).wait(); + + // Now, run the kernel by first getting its image as an executable, + // making an L0 kernel out of it and then making a SYCL kernel out of + // the L0 kernel. Run this kernel on the SYCL API and verify + // that it has the same result as the kernel that was run directly on SYCL API. + // First, get a kernel bundle that contains the kernel "iota". + auto exe_bndl = sycl::get_kernel_bundle( + ctxt, {d}, + [&](const sycl::device_image &img) { + return img.has_kernel(iota_id, d); + }); + std::vector bytes; + const sycl::device_image &img = + *(exe_bndl.begin()); + bytes = img.ext_oneapi_get_backend_content(); + + auto ZeContext = sycl::get_native(ctxt); + auto ZeDevice = sycl::get_native(d); + + ze_result_t status; + ze_module_desc_t moduleDesc = { + ZE_STRUCTURE_TYPE_MODULE_DESC, + nullptr, + ZE_MODULE_FORMAT_IL_SPIRV, + bytes.size(), + reinterpret_cast(bytes.data()), + nullptr, + nullptr}; + ze_module_handle_t ZeModule; + status = zeModuleCreate(ZeContext, ZeDevice, &moduleDesc, &ZeModule, nullptr); + assert(status == ZE_RESULT_SUCCESS); + + ze_kernel_desc_t kernelDesc = {ZE_STRUCTURE_TYPE_KERNEL_DESC, nullptr, 0, + "__sycl_kernel_iota"}; + ze_kernel_handle_t ZeKernel; + status = zeKernelCreate(ZeModule, &kernelDesc, &ZeKernel); + assert(status == ZE_RESULT_SUCCESS); + sycl::kernel k_iota_twin = + sycl::make_kernel( + {sycl::make_kernel_bundle({ZeModule}, ctxt), ZeKernel}, ctxt); + int *ptr_twin = sycl::malloc_shared(1, q); + q.submit([&](sycl::handler &cgh) { + cgh.set_args(ptr_twin); + cgh.parallel_for(sycl::nd_range{{1}, {1}}, k_iota_twin); + }).wait(); + assert(*ptr_twin == *ptr); + sycl::free(ptr, q); + sycl::free(ptr_twin, q); +#endif +} From 12bdf3f56d825daa6f80df4851cbdda134ca2a54 Mon Sep 17 00:00:00 2001 From: "Bushi, Lorenc" Date: Wed, 15 Jan 2025 13:10:14 -0800 Subject: [PATCH 03/53] Change extension spec status section --- ...cl_ext_oneapi_device_image_backend_content.asciidoc | 10 ++++------ 1 file changed, 4 insertions(+), 6 deletions(-) diff --git a/sycl/doc/extensions/experimental/sycl_ext_oneapi_device_image_backend_content.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_oneapi_device_image_backend_content.asciidoc index 566e0f2482334..2f581e220873a 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_device_image_backend_content.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_device_image_backend_content.asciidoc @@ -44,12 +44,10 @@ the SYCL specification refer to that revision. == Status -This is a proposed extension specification, intended to gather community -feedback. -Interfaces defined in this specification may not be implemented yet or may be -in a preliminary state. -The specification itself may also change in incompatible ways before it is -finalized. +This is an experimental extension specification, intended to provide early +access to features and gather community feedback. Interfaces defined in this +specification are implemented in {dpcpp}, but they are not finalized and may +change incompatibly in future versions of {dpcpp} without prior notice. *Shipping software products should not rely on APIs defined in this specification.* From c15a31a77ec869af414a9faeb7b05fdbc60b6e63 Mon Sep 17 00:00:00 2001 From: Lorenc Bushi Date: Thu, 16 Jan 2025 12:33:01 -0800 Subject: [PATCH 04/53] Add more tests and fix ABI related failures --- sycl/include/sycl/kernel_bundle.hpp | 8 ++- .../DeviceImageBackendContent/basic_test.cpp | 72 +++++++++++++++++++ sycl/test/abi/sycl_symbols_linux.dump | 5 +- 3 files changed, 82 insertions(+), 3 deletions(-) create mode 100644 sycl/test-e2e/DeviceImageBackendContent/basic_test.cpp diff --git a/sycl/include/sycl/kernel_bundle.hpp b/sycl/include/sycl/kernel_bundle.hpp index 12dda72c1b45c..02c6be556a68d 100644 --- a/sycl/include/sycl/kernel_bundle.hpp +++ b/sycl/include/sycl/kernel_bundle.hpp @@ -32,12 +32,12 @@ #include // for function #include // for distance #include // for shared_ptr, operator==, hash +#include // for span #include // for string #include // for enable_if_t, remove_refer... #include // for move #include // for hash #include // for vector -#include // for span namespace sycl { inline namespace _V1 { @@ -120,9 +120,11 @@ class __SYCL_EXPORT device_image_plain { backend get_backend() const; +#if (!defined(_HAS_STD_BYTE) || _HAS_STD_BYTE != 0) const std::byte *get_BinaryStart() const; const std::byte *get_BinaryEnd() const; +#endif template friend const decltype(Obj::impl) & @@ -156,6 +158,7 @@ class device_image : public detail::device_image_plain, return device_image_plain::get_backend(); } +#if (!defined(_HAS_STD_BYTE) || _HAS_STD_BYTE != 0) template > std::vector ext_oneapi_get_backend_content() const { @@ -170,7 +173,8 @@ class device_image : public detail::device_image_plain, return std::span(device_image_plain::get_BinaryStart(), device_image_plain::get_BinaryEnd()); } -#endif +#endif // __cpp_lib_span +#endif // _HAS_STD_BYTE private: device_image(detail::DeviceImageImplPtr Impl) diff --git a/sycl/test-e2e/DeviceImageBackendContent/basic_test.cpp b/sycl/test-e2e/DeviceImageBackendContent/basic_test.cpp new file mode 100644 index 0000000000000..69dac913418ca --- /dev/null +++ b/sycl/test-e2e/DeviceImageBackendContent/basic_test.cpp @@ -0,0 +1,72 @@ +// RUN: %{build} -fsyntax-only -DTEST_API_VIOLATION=1 -Xclang -verify -Xclang -verify-ignore-unexpected=note,warning +// RUN: %{build} -o %t.out +// RUN: %{run} %t.out + +#ifdef __cpp_lib_span +#include +#endif +#include +#include + +class kernel; + +void define_kernel(sycl::queue &q) { + int data; + sycl::buffer data_buf(&data, 1); + q.submit([&](sycl::handler &cgh) { + sycl::accessor data_acc(data_buf, cgh); + cgh.parallel_for( + sycl::nd_range{{1}, {1}}, + [=](sycl::nd_item<> it) { data_acc[0] = 42; }); + }); +} + +int main() { + sycl::device d; + sycl::queue q{d}; + sycl::context ctxt = q.get_context(); + sycl::kernel_id id = sycl::get_kernel_id(); + auto bundle = + sycl::get_kernel_bundle(ctxt, {id}); + assert(!bundle.empty()); + sycl::backend backend; + std::vector bytes; +#ifdef __cpp_lib_span + std::span bytes_view; +#endif + for (const auto &img : bundle) { + if (img.has_kernel(id, d)) { + // Check that all 3 functions of the api compile. + // Furthermore, check that the backend corresponds to the backend of the + // bundle Check that the view of the content is indeed equal to the + // content. + backend = img.ext_oneapi_get_backend(); + assert(backend == bundle.get_backend()); + bytes = img.ext_oneapi_get_backend_content(); +#ifdef __cpp_lib_span + bytes_view = img.ext_oneapi_get_backend_content_view(); + assert(bytes_view.size() == bytes.size()); + for (size_t i = 0; i < bytes.size(); ++i) { + assert(bytes[i] == bytes_view[i]); + } +#endif + } + } + +#ifdef TEST_API_VIOLATION + // Check that the ext_oneapi_get_backend_content and the + // ext_oneapi_get_backend_content_view of the content functions are only + // available + // when the image is in the executable state. + + auto input_bundle = + sycl::get_kernel_bundle(ctxt, {id}); + // expected-error@+1 {{no matching member function for call to 'ext_oneapi_get_backend_content'}} + bytes = (*input_bundle.begin()).ext_oneapi_get_backend_content(); +#ifdef _cpp_lib_span + // expected-error@+1 {{no matching member function for call to 'ext_oneapi_get_backend_content_view'}} + bytes_view = (*input_bundle.begin()).ext_oneapi_get_backend_content_view(); +#endif // __cpp_lib_span +#endif // TEST_API_VIOLATION + return 0; +} diff --git a/sycl/test/abi/sycl_symbols_linux.dump b/sycl/test/abi/sycl_symbols_linux.dump index 26a129e33ef85..05b78e77c6521 100644 --- a/sycl/test/abi/sycl_symbols_linux.dump +++ b/sycl/test/abi/sycl_symbols_linux.dump @@ -3524,6 +3524,7 @@ _ZN4sycl3_V17handler20associateWithHandlerEPNS0_6detail30UnsampledImageAccessorB _ZN4sycl3_V17handler20memcpyToDeviceGlobalEPKvS3_bmm _ZN4sycl3_V17handler20setKernelCacheConfigENS1_23StableKernelCacheConfigE _ZN4sycl3_V17handler20setStateSpecConstSetEv +_ZN4sycl3_V17handler21setKernelWorkGroupMemEm _ZN4sycl3_V17handler21setUserFacingNodeTypeENS0_3ext6oneapi12experimental9node_typeE _ZN4sycl3_V17handler22ext_oneapi_fill2d_implEPvmPKvmmm _ZN4sycl3_V17handler22memcpyFromDeviceGlobalEPvPKvbmm @@ -3531,7 +3532,6 @@ _ZN4sycl3_V17handler22setHandlerKernelBundleENS0_6kernelE _ZN4sycl3_V17handler22setHandlerKernelBundleERKSt10shared_ptrINS0_6detail18kernel_bundle_implEE _ZN4sycl3_V17handler22setKernelClusterLaunchENS0_5rangeILi3EEEi _ZN4sycl3_V17handler22setKernelIsCooperativeEb -_ZN4sycl3_V17handler21setKernelWorkGroupMemEm _ZN4sycl3_V17handler24GetRangeRoundingSettingsERmS2_S2_ _ZN4sycl3_V17handler24ext_intel_read_host_pipeENS0_6detail11string_viewEPvmb _ZN4sycl3_V17handler24ext_oneapi_memcpy2d_implEPvmPKvmmm @@ -3737,6 +3737,9 @@ _ZNK4sycl3_V16detail16AccessorBaseHost6getPtrEv _ZNK4sycl3_V16detail16AccessorBaseHost9getOffsetEv _ZNK4sycl3_V16detail18device_image_plain10has_kernelERKNS0_9kernel_idE _ZNK4sycl3_V16detail18device_image_plain10has_kernelERKNS0_9kernel_idERKNS0_6deviceE +_ZNK4sycl3_V16detail18device_image_plain11get_backendEv +_ZNK4sycl3_V16detail18device_image_plain13get_BinaryEndEv +_ZNK4sycl3_V16detail18device_image_plain15get_BinaryStartEv _ZNK4sycl3_V16detail18device_image_plain9getNativeEv _ZNK4sycl3_V16detail19kernel_bundle_plain10get_kernelERKNS0_9kernel_idE _ZNK4sycl3_V16detail19kernel_bundle_plain10has_kernelERKNS0_9kernel_idE From 0ba58c7cc7f3e32cb07d2fb95e7fedb8ce58bc2a Mon Sep 17 00:00:00 2001 From: Lorenc Bushi Date: Thu, 16 Jan 2025 12:35:59 -0800 Subject: [PATCH 05/53] Add checks for span feature --- sycl/include/sycl/kernel_bundle.hpp | 2 ++ 1 file changed, 2 insertions(+) diff --git a/sycl/include/sycl/kernel_bundle.hpp b/sycl/include/sycl/kernel_bundle.hpp index 02c6be556a68d..c442b169c4c72 100644 --- a/sycl/include/sycl/kernel_bundle.hpp +++ b/sycl/include/sycl/kernel_bundle.hpp @@ -32,7 +32,9 @@ #include // for function #include // for distance #include // for shared_ptr, operator==, hash +#ifdef __cpp_lib_span #include // for span +#endif #include // for string #include // for enable_if_t, remove_refer... #include // for move From 50018cbe17545b620c008512c6db40858f169bfd Mon Sep 17 00:00:00 2001 From: Lorenc Bushi Date: Thu, 16 Jan 2025 18:46:20 -0800 Subject: [PATCH 06/53] Refactoring --- sycl/include/sycl/kernel_bundle.hpp | 25 +++++++++++++------------ sycl/source/kernel_bundle.cpp | 20 +++++++++++++------- 2 files changed, 26 insertions(+), 19 deletions(-) diff --git a/sycl/include/sycl/kernel_bundle.hpp b/sycl/include/sycl/kernel_bundle.hpp index c442b169c4c72..cc012e116822a 100644 --- a/sycl/include/sycl/kernel_bundle.hpp +++ b/sycl/include/sycl/kernel_bundle.hpp @@ -117,16 +117,19 @@ class __SYCL_EXPORT device_image_plain { ur_native_handle_t getNative() const; -protected: - detail::DeviceImageImplPtr impl; - - backend get_backend() const; + backend ext_oneapi_get_backend() const noexcept; #if (!defined(_HAS_STD_BYTE) || _HAS_STD_BYTE != 0) - const std::byte *get_BinaryStart() const; + std::vector ext_oneapi_get_backend_content() const; - const std::byte *get_BinaryEnd() const; -#endif +#ifdef __cpp_lib_span + std::span ext_oneapi_get_backend_content_view() const; +#endif // __cpp_lib_span + +#endif // HAS_STD_BYTE + +protected: + detail::DeviceImageImplPtr impl; template friend const decltype(Obj::impl) & @@ -157,23 +160,21 @@ class device_image : public detail::device_image_plain, } backend ext_oneapi_get_backend() const noexcept { - return device_image_plain::get_backend(); + return device_image_plain::ext_oneapi_get_backend(); } #if (!defined(_HAS_STD_BYTE) || _HAS_STD_BYTE != 0) template > std::vector ext_oneapi_get_backend_content() const { - return std::vector(device_image_plain::get_BinaryStart(), - device_image_plain::get_BinaryEnd()); + return device_image_plain::ext_oneapi_get_backend_content(); } #ifdef __cpp_lib_span template > std::span ext_oneapi_get_content_backend_view() const { - return std::span(device_image_plain::get_BinaryStart(), - device_image_plain::get_BinaryEnd()); + return device_image_plain::ext_oneapi_get_backend_content_view(); } #endif // __cpp_lib_span #endif // _HAS_STD_BYTE diff --git a/sycl/source/kernel_bundle.cpp b/sycl/source/kernel_bundle.cpp index bfc307c6bb17f..f6a8818152c19 100644 --- a/sycl/source/kernel_bundle.cpp +++ b/sycl/source/kernel_bundle.cpp @@ -45,19 +45,25 @@ ur_native_handle_t device_image_plain::getNative() const { return impl->getNative(); } -backend device_image_plain::get_backend() const { +backend device_image_plain::ext_oneapi_get_backend() const noexcept{ return impl->get_context().get_backend(); } -const std::byte *device_image_plain::get_BinaryStart() const { - return reinterpret_cast( - impl->get_bin_image_ref()->getRawData().BinaryStart); +std::vector device_image_plain::ext_oneapi_get_backend_content() const { + return std::vector(reinterpret_cast( + impl->get_bin_image_ref()->getRawData().BinaryStart), + reinterpret_cast( + impl->get_bin_image_ref()->getRawData().BinaryEnd)); } -const std::byte *device_image_plain::get_BinaryEnd() const { - return reinterpret_cast( - impl->get_bin_image_ref()->getRawData().BinaryEnd); +#ifdef __cpp_lib_span +std::span device_image_plain::ext_oneapi_get_backend_content_view() const { + return std::span(reinterpret_cast( + impl->get_bin_image_ref()->getRawData().BinaryStart), + reinterpret_cast( + impl->get_bin_image_ref()->getRawData().BinaryEnd)); } +#endif //////////////////////////// ///// kernel_bundle_plain From b20f04219e190f34587bfae4defa12c7af556cd3 Mon Sep 17 00:00:00 2001 From: Lorenc Bushi Date: Thu, 16 Jan 2025 19:48:33 -0800 Subject: [PATCH 07/53] More refactoring --- sycl/include/sycl/kernel_bundle.hpp | 14 +++++------ .../L0_interop_test.cpp | 24 +++++++++++-------- sycl/test/abi/sycl_symbols_linux.dump | 5 ++-- 3 files changed, 23 insertions(+), 20 deletions(-) diff --git a/sycl/include/sycl/kernel_bundle.hpp b/sycl/include/sycl/kernel_bundle.hpp index cc012e116822a..c3060c7c96f73 100644 --- a/sycl/include/sycl/kernel_bundle.hpp +++ b/sycl/include/sycl/kernel_bundle.hpp @@ -26,14 +26,14 @@ #include // build_options #include // and log -#include // for array -#include // for std::byte -#include // for size_t, memcpy -#include // for function -#include // for distance -#include // for shared_ptr, operator==, hash +#include // for array +#include // for std::byte +#include // for size_t, memcpy +#include // for function +#include // for distance +#include // for shared_ptr, operator==, hash #ifdef __cpp_lib_span -#include // for span +#include // for span #endif #include // for string #include // for enable_if_t, remove_refer... diff --git a/sycl/test-e2e/DeviceImageBackendContent/L0_interop_test.cpp b/sycl/test-e2e/DeviceImageBackendContent/L0_interop_test.cpp index aa77862f08429..06a764f0ac690 100644 --- a/sycl/test-e2e/DeviceImageBackendContent/L0_interop_test.cpp +++ b/sycl/test-e2e/DeviceImageBackendContent/L0_interop_test.cpp @@ -3,8 +3,8 @@ // RUN: %{run} %t.out // #include -#include #include +#include #include #include #include @@ -12,8 +12,8 @@ namespace syclext = sycl::ext::oneapi; namespace syclexp = sycl::ext::oneapi::experimental; -extern"C" SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<1>)) -void iota(int *ptr) { +extern "C" SYCL_EXT_ONEAPI_FUNCTION_PROPERTY( + (syclexp::nd_range_kernel<1>)) void iota(int *ptr) { size_t id = syclext::this_work_item::get_nd_item<1>().get_global_linear_id(); ptr[id] = 42; } @@ -36,15 +36,15 @@ int main() { cgh.parallel_for(sycl::nd_range{{1}, {1}}, k_iota); }).wait(); - // Now, run the kernel by first getting its image as an executable, - // making an L0 kernel out of it and then making a SYCL kernel out of - // the L0 kernel. Run this kernel on the SYCL API and verify - // that it has the same result as the kernel that was run directly on SYCL API. - // First, get a kernel bundle that contains the kernel "iota". + // Now, run the kernel by first getting its image as an executable, + // making an L0 kernel out of it and then making a SYCL kernel out of + // the L0 kernel. Run this kernel on the SYCL API and verify + // that it has the same result as the kernel that was run directly on SYCL + // API. First, get a kernel bundle that contains the kernel "iota". auto exe_bndl = sycl::get_kernel_bundle( ctxt, {d}, [&](const sycl::device_image &img) { - return img.has_kernel(iota_id, d); + return img.has_kernel(iota_id, d); }); std::vector bytes; const sycl::device_image &img = @@ -74,7 +74,11 @@ int main() { assert(status == ZE_RESULT_SUCCESS); sycl::kernel k_iota_twin = sycl::make_kernel( - {sycl::make_kernel_bundle({ZeModule}, ctxt), ZeKernel}, ctxt); + {sycl::make_kernel_bundle({ZeModule}, + ctxt), + ZeKernel}, + ctxt); int *ptr_twin = sycl::malloc_shared(1, q); q.submit([&](sycl::handler &cgh) { cgh.set_args(ptr_twin); diff --git a/sycl/test/abi/sycl_symbols_linux.dump b/sycl/test/abi/sycl_symbols_linux.dump index 05b78e77c6521..c808f03f1fa4a 100644 --- a/sycl/test/abi/sycl_symbols_linux.dump +++ b/sycl/test/abi/sycl_symbols_linux.dump @@ -3737,9 +3737,8 @@ _ZNK4sycl3_V16detail16AccessorBaseHost6getPtrEv _ZNK4sycl3_V16detail16AccessorBaseHost9getOffsetEv _ZNK4sycl3_V16detail18device_image_plain10has_kernelERKNS0_9kernel_idE _ZNK4sycl3_V16detail18device_image_plain10has_kernelERKNS0_9kernel_idERKNS0_6deviceE -_ZNK4sycl3_V16detail18device_image_plain11get_backendEv -_ZNK4sycl3_V16detail18device_image_plain13get_BinaryEndEv -_ZNK4sycl3_V16detail18device_image_plain15get_BinaryStartEv +_ZNK4sycl3_V16detail18device_image_plain22ext_oneapi_get_backendEv +_ZNK4sycl3_V16detail18device_image_plain30ext_oneapi_get_backend_contentEv _ZNK4sycl3_V16detail18device_image_plain9getNativeEv _ZNK4sycl3_V16detail19kernel_bundle_plain10get_kernelERKNS0_9kernel_idE _ZNK4sycl3_V16detail19kernel_bundle_plain10has_kernelERKNS0_9kernel_idE From 2c7b110bdabc37ddb31767136e7042e5c16b2314 Mon Sep 17 00:00:00 2001 From: Lorenc Bushi Date: Thu, 16 Jan 2025 20:17:44 -0800 Subject: [PATCH 08/53] Formatting --- sycl/source/kernel_bundle.cpp | 10 +++++----- 1 file changed, 5 insertions(+), 5 deletions(-) diff --git a/sycl/source/kernel_bundle.cpp b/sycl/source/kernel_bundle.cpp index f6a8818152c19..1cbbe0e190f69 100644 --- a/sycl/source/kernel_bundle.cpp +++ b/sycl/source/kernel_bundle.cpp @@ -45,25 +45,25 @@ ur_native_handle_t device_image_plain::getNative() const { return impl->getNative(); } -backend device_image_plain::ext_oneapi_get_backend() const noexcept{ +backend device_image_plain::ext_oneapi_get_backend() const noexcept { return impl->get_context().get_backend(); } -std::vector device_image_plain::ext_oneapi_get_backend_content() const { +std::vector +device_image_plain::ext_oneapi_get_backend_content() const { return std::vector(reinterpret_cast( impl->get_bin_image_ref()->getRawData().BinaryStart), reinterpret_cast( impl->get_bin_image_ref()->getRawData().BinaryEnd)); } -#ifdef __cpp_lib_span -std::span device_image_plain::ext_oneapi_get_backend_content_view() const { +std::span +device_image_plain::ext_oneapi_get_backend_content_view() const { return std::span(reinterpret_cast( impl->get_bin_image_ref()->getRawData().BinaryStart), reinterpret_cast( impl->get_bin_image_ref()->getRawData().BinaryEnd)); } -#endif //////////////////////////// ///// kernel_bundle_plain From 36f4095014526bb4843af381b35729a0d35bcfeb Mon Sep 17 00:00:00 2001 From: "Bushi, Lorenc" Date: Fri, 17 Jan 2025 09:57:59 -0800 Subject: [PATCH 09/53] Add windows symbols --- sycl/test/abi/sycl_symbols_windows.dump | 28 +++++++++++++------------ 1 file changed, 15 insertions(+), 13 deletions(-) diff --git a/sycl/test/abi/sycl_symbols_windows.dump b/sycl/test/abi/sycl_symbols_windows.dump index 1519f0b0bf3a7..5d09568968d70 100644 --- a/sycl/test/abi/sycl_symbols_windows.dump +++ b/sycl/test/abi/sycl_symbols_windows.dump @@ -282,9 +282,9 @@ ??0SampledImageAccessorBaseHost@detail@_V1@sycl@@QEAA@$$QEAV0123@@Z ??0SampledImageAccessorBaseHost@detail@_V1@sycl@@QEAA@AEBV0123@@Z ??0SampledImageAccessorBaseHost@detail@_V1@sycl@@QEAA@V?$range@$02@23@PEAXHHV?$id@$02@23@W4image_channel_type@23@W4image_channel_order@23@Uimage_sampler@23@AEBVproperty_list@23@@Z -??0SubmissionInfo@detail@_V1@sycl@@QEAA@XZ -??0SubmissionInfo@detail@_V1@sycl@@QEAA@AEBV0123@@Z ??0SubmissionInfo@detail@_V1@sycl@@QEAA@$$QEAV0123@@Z +??0SubmissionInfo@detail@_V1@sycl@@QEAA@AEBV0123@@Z +??0SubmissionInfo@detail@_V1@sycl@@QEAA@XZ ??0UnsampledImageAccessorBaseHost@detail@_V1@sycl@@IEAA@AEBV?$shared_ptr@VUnsampledImageAccessorImplHost@detail@_V1@sycl@@@std@@@Z ??0UnsampledImageAccessorBaseHost@detail@_V1@sycl@@QEAA@$$QEAV0123@@Z ??0UnsampledImageAccessorBaseHost@detail@_V1@sycl@@QEAA@AEBV0123@@Z @@ -330,18 +330,12 @@ ??0device_image_plain@detail@_V1@sycl@@QEAA@AEBV?$shared_ptr@Vdevice_image_impl@detail@_V1@sycl@@@std@@@Z ??0device_selector@_V1@sycl@@QEAA@AEBV012@@Z ??0device_selector@_V1@sycl@@QEAA@XZ +??0dynamic_command_group@experimental@oneapi@ext@_V1@sycl@@QEAA@$$QEAV012345@@Z +??0dynamic_command_group@experimental@oneapi@ext@_V1@sycl@@QEAA@AEBV012345@@Z +??0dynamic_command_group@experimental@oneapi@ext@_V1@sycl@@QEAA@AEBV?$command_graph@$0A@@12345@AEBV?$vector@V?$function@$$A6AXAEAVhandler@_V1@sycl@@@Z@std@@V?$allocator@V?$function@$$A6AXAEAVhandler@_V1@sycl@@@Z@std@@@2@@std@@@Z ??0dynamic_parameter_base@detail@experimental@oneapi@ext@_V1@sycl@@QEAA@$$QEAV0123456@@Z ??0dynamic_parameter_base@detail@experimental@oneapi@ext@_V1@sycl@@QEAA@AEBV0123456@@Z ??0dynamic_parameter_base@detail@experimental@oneapi@ext@_V1@sycl@@QEAA@V?$command_graph@$0A@@23456@_KPEBX@Z -?get_active_index@dynamic_command_group@experimental@oneapi@ext@_V1@sycl@@QEBA_KXZ -??4dynamic_command_group@experimental@oneapi@ext@_V1@sycl@@QEAAAEAV012345@$$QEAV012345@@Z -?addImpl@modifiable_command_graph@detail@experimental@oneapi@ext@_V1@sycl@@IEAA?AVnode@34567@AEAVdynamic_command_group@34567@AEBV?$vector@Vnode@experimental@oneapi@ext@_V1@sycl@@V?$allocator@Vnode@experimental@oneapi@ext@_V1@sycl@@@std@@@std@@@Z -??0dynamic_command_group@experimental@oneapi@ext@_V1@sycl@@QEAA@AEBV?$command_graph@$0A@@12345@AEBV?$vector@V?$function@$$A6AXAEAVhandler@_V1@sycl@@@Z@std@@V?$allocator@V?$function@$$A6AXAEAVhandler@_V1@sycl@@@Z@std@@@2@@std@@@Z -??0dynamic_command_group@experimental@oneapi@ext@_V1@sycl@@QEAA@$$QEAV012345@@Z -??0dynamic_command_group@experimental@oneapi@ext@_V1@sycl@@QEAA@AEBV012345@@Z -?set_active_index@dynamic_command_group@experimental@oneapi@ext@_V1@sycl@@QEAAX_K@Z -??1dynamic_command_group@experimental@oneapi@ext@_V1@sycl@@QEAA@XZ -??4dynamic_command_group@experimental@oneapi@ext@_V1@sycl@@QEAAAEAV012345@AEBV012345@@Z ??0event@_V1@sycl@@AEAA@V?$shared_ptr@Vevent_impl@detail@_V1@sycl@@@std@@@Z ??0event@_V1@sycl@@QEAA@$$QEAV012@@Z ??0event@_V1@sycl@@QEAA@AEBV012@@Z @@ -479,6 +473,7 @@ ??1device@_V1@sycl@@QEAA@XZ ??1device_image_plain@detail@_V1@sycl@@QEAA@XZ ??1device_selector@_V1@sycl@@UEAA@XZ +??1dynamic_command_group@experimental@oneapi@ext@_V1@sycl@@QEAA@XZ ??1dynamic_parameter_base@detail@experimental@oneapi@ext@_V1@sycl@@QEAA@XZ ??1event@_V1@sycl@@QEAA@XZ ??1exception@_V1@sycl@@UEAA@XZ @@ -556,6 +551,8 @@ ??4device_image_plain@detail@_V1@sycl@@QEAAAEAV0123@$$QEAV0123@@Z ??4device_image_plain@detail@_V1@sycl@@QEAAAEAV0123@AEBV0123@@Z ??4device_selector@_V1@sycl@@QEAAAEAV012@AEBV012@@Z +??4dynamic_command_group@experimental@oneapi@ext@_V1@sycl@@QEAAAEAV012345@$$QEAV012345@@Z +??4dynamic_command_group@experimental@oneapi@ext@_V1@sycl@@QEAAAEAV012345@AEBV012345@@Z ??4dynamic_parameter_base@detail@experimental@oneapi@ext@_V1@sycl@@QEAAAEAV0123456@$$QEAV0123456@@Z ??4dynamic_parameter_base@detail@experimental@oneapi@ext@_V1@sycl@@QEAAAEAV0123456@AEBV0123456@@Z ??4event@_V1@sycl@@QEAAAEAV012@$$QEAV012@@Z @@ -654,8 +651,8 @@ ?GDBMethodsAnchor@UnsampledImageAccessorBaseHost@detail@_V1@sycl@@IEAAXXZ ?GetRangeRoundingSettings@handler@_V1@sycl@@AEAAXAEA_K00@Z ?HasAssociatedAccessor@handler@_V1@sycl@@AEBA_NPEAVAccessorImplHost@detail@23@W4target@access@23@@Z -?PostProcessorFunc@SubmissionInfo@detail@_V1@sycl@@QEBAAEBV?$optional@V?$function@$$A6AX_N0AEAVevent@_V1@sycl@@@Z@std@@@234@XZ ?PostProcessorFunc@SubmissionInfo@detail@_V1@sycl@@QEAAAEAV?$optional@V?$function@$$A6AX_N0AEAVevent@_V1@sycl@@@Z@std@@@234@XZ +?PostProcessorFunc@SubmissionInfo@detail@_V1@sycl@@QEBAAEBV?$optional@V?$function@$$A6AX_N0AEAVevent@_V1@sycl@@@Z@std@@@234@XZ ?PushBack@exception_list@_V1@sycl@@AEAAX$$QEAVexception_ptr@std@@@Z ?PushBack@exception_list@_V1@sycl@@AEAAXAEBVexception_ptr@std@@@Z ?RangeRoundingTrace@handler@_V1@sycl@@AEAA_NXZ @@ -3707,6 +3704,7 @@ ?addHostAccessorAndWait@detail@_V1@sycl@@YAXPEAVAccessorImplHost@123@@Z ?addHostSampledImageAccessorAndWait@detail@_V1@sycl@@YAXPEAVSampledImageAccessorImplHost@123@@Z ?addHostUnsampledImageAccessorAndWait@detail@_V1@sycl@@YAXPEAVUnsampledImageAccessorImplHost@123@@Z +?addImpl@modifiable_command_graph@detail@experimental@oneapi@ext@_V1@sycl@@IEAA?AVnode@34567@AEAVdynamic_command_group@34567@AEBV?$vector@Vnode@experimental@oneapi@ext@_V1@sycl@@V?$allocator@Vnode@experimental@oneapi@ext@_V1@sycl@@@std@@@std@@@Z ?addImpl@modifiable_command_graph@detail@experimental@oneapi@ext@_V1@sycl@@IEAA?AVnode@34567@AEBV?$vector@Vnode@experimental@oneapi@ext@_V1@sycl@@V?$allocator@Vnode@experimental@oneapi@ext@_V1@sycl@@@std@@@std@@@Z ?addImpl@modifiable_command_graph@detail@experimental@oneapi@ext@_V1@sycl@@IEAA?AVnode@34567@V?$function@$$A6AXAEAVhandler@_V1@sycl@@@Z@std@@AEBV?$vector@Vnode@experimental@oneapi@ext@_V1@sycl@@V?$allocator@Vnode@experimental@oneapi@ext@_V1@sycl@@@std@@@std@@@Z ?addLifetimeSharedPtrStorage@handler@_V1@sycl@@AEAAXV?$shared_ptr@$$CBX@std@@@Z @@ -3842,13 +3840,15 @@ ?ext_oneapi_empty@queue@_V1@sycl@@QEBA_NXZ ?ext_oneapi_enable_peer_access@device@_V1@sycl@@QEAAXAEBV123@@Z ?ext_oneapi_fill2d_impl@handler@_V1@sycl@@AEAAXPEAX_KPEBX111@Z +?ext_oneapi_get_backend@device_image_plain@detail@_V1@sycl@@QEBA?AW4backend@34@XZ +?ext_oneapi_get_backend_content@device_image_plain@detail@_V1@sycl@@QEBA?AV?$vector@W4byte@std@@V?$allocator@W4byte@std@@@2@@std@@XZ ?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 ?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 -?ext_oneapi_get_last_event_impl@queue@_V1@sycl@@AEBA?AV?$optional@Vevent@_V1@sycl@@@detail@23@XZ ?ext_oneapi_get_last_event@queue@_V1@sycl@@QEBA?AV?$optional@Vevent@_V1@sycl@@@std@@XZ +?ext_oneapi_get_last_event_impl@queue@_V1@sycl@@AEBA?AV?$optional@Vevent@_V1@sycl@@@detail@23@XZ ?ext_oneapi_get_state@queue@_V1@sycl@@QEBA?AW4queue_state@experimental@oneapi@ext@23@XZ ?ext_oneapi_graph@handler@_V1@sycl@@QEAAXV?$command_graph@$00@experimental@oneapi@ext@23@@Z ?ext_oneapi_graph@queue@_V1@sycl@@QEAA?AVevent@23@V?$command_graph@$00@experimental@oneapi@ext@23@AEBUcode_location@detail@23@@Z @@ -4014,6 +4014,7 @@ ?getType@handler@_V1@sycl@@AEBA?AW4CGType@detail@23@XZ ?getValueFromDynamicParameter@detail@_V1@sycl@@YAPEAXAEAVdynamic_parameter_base@1experimental@oneapi@ext@23@@Z ?get_access_mode@experimental@oneapi@ext@_V1@sycl@@YA?AW4address_access_mode@12345@PEBX_KAEBVcontext@45@@Z +?get_active_index@dynamic_command_group@experimental@oneapi@ext@_V1@sycl@@QEBA_KXZ ?get_addressing_mode@sampler@_V1@sycl@@QEBA?AW4addressing_mode@23@XZ ?get_allocator_internal@buffer_plain@detail@_V1@sycl@@IEBAAEBV?$unique_ptr@VSYCLMemObjAllocator@detail@_V1@sycl@@U?$default_delete@VSYCLMemObjAllocator@detail@_V1@sycl@@@std@@@std@@XZ ?get_allocator_internal@image_plain@detail@_V1@sycl@@IEBAAEBV?$unique_ptr@VSYCLMemObjAllocator@detail@_V1@sycl@@U?$default_delete@VSYCLMemObjAllocator@detail@_V1@sycl@@@std@@@std@@XZ @@ -4274,6 +4275,7 @@ ?setType@handler@_V1@sycl@@AEAAXW4CGType@detail@23@@Z ?setUserFacingNodeType@handler@_V1@sycl@@AEAAXW4node_type@experimental@oneapi@ext@23@@Z ?set_access_mode@experimental@oneapi@ext@_V1@sycl@@YAXPEBX_KW4address_access_mode@12345@AEBVcontext@45@@Z +?set_active_index@dynamic_command_group@experimental@oneapi@ext@_V1@sycl@@QEAAX_K@Z ?set_arg@handler@_V1@sycl@@QEAAXH$$QEAVraw_kernel_arg@experimental@oneapi@ext@23@@Z ?set_final_data_internal@buffer_plain@detail@_V1@sycl@@IEAAXAEBV?$function@$$A6AXAEBV?$function@$$A6AXPEAX@Z@std@@@Z@std@@@Z ?set_final_data_internal@buffer_plain@detail@_V1@sycl@@IEAAXXZ From 5d4a6bf5893dcdbde5c89c94f02bd8f9c4fcbf85 Mon Sep 17 00:00:00 2001 From: "Bushi, Lorenc" Date: Fri, 17 Jan 2025 10:01:58 -0800 Subject: [PATCH 10/53] Add check for std::span feature --- sycl/source/kernel_bundle.cpp | 2 ++ 1 file changed, 2 insertions(+) diff --git a/sycl/source/kernel_bundle.cpp b/sycl/source/kernel_bundle.cpp index 1cbbe0e190f69..3e987de513bb4 100644 --- a/sycl/source/kernel_bundle.cpp +++ b/sycl/source/kernel_bundle.cpp @@ -57,6 +57,7 @@ device_image_plain::ext_oneapi_get_backend_content() const { impl->get_bin_image_ref()->getRawData().BinaryEnd)); } +#ifdef __cpp_lib_span std::span device_image_plain::ext_oneapi_get_backend_content_view() const { return std::span(reinterpret_cast( @@ -64,6 +65,7 @@ device_image_plain::ext_oneapi_get_backend_content_view() const { reinterpret_cast( impl->get_bin_image_ref()->getRawData().BinaryEnd)); } +#endif //////////////////////////// ///// kernel_bundle_plain From 273f6df374a4e1091b8f0da64e21ce788949c836 Mon Sep 17 00:00:00 2001 From: "Bushi, Lorenc" Date: Fri, 17 Jan 2025 10:09:38 -0800 Subject: [PATCH 11/53] Add comments explaining limitations of free function kernel usage --- sycl/test-e2e/DeviceImageBackendContent/L0_interop_test.cpp | 3 +++ 1 file changed, 3 insertions(+) diff --git a/sycl/test-e2e/DeviceImageBackendContent/L0_interop_test.cpp b/sycl/test-e2e/DeviceImageBackendContent/L0_interop_test.cpp index 06a764f0ac690..21b52a9fab30c 100644 --- a/sycl/test-e2e/DeviceImageBackendContent/L0_interop_test.cpp +++ b/sycl/test-e2e/DeviceImageBackendContent/L0_interop_test.cpp @@ -25,6 +25,9 @@ int main() { sycl::queue q{d}; sycl::context ctxt = q.get_context(); + // The following ifndef is required due to a number of limitations of free + // function kernels. See CMPLRLLVM-61498. + // TODO: Remove it once these limitations are no longer there. #ifndef __SYCL_DEVICE_ONLY__ // First, run the kernel using the SYCL API. auto Bundle = sycl::get_kernel_bundle(ctxt); From 192cf75fdaeaef98b64c6e63a2b2ea2a008990a3 Mon Sep 17 00:00:00 2001 From: "Bushi, Lorenc" Date: Fri, 17 Jan 2025 10:13:01 -0800 Subject: [PATCH 12/53] Add asserts for extra safety --- sycl/test-e2e/DeviceImageBackendContent/L0_interop_test.cpp | 1 + 1 file changed, 1 insertion(+) diff --git a/sycl/test-e2e/DeviceImageBackendContent/L0_interop_test.cpp b/sycl/test-e2e/DeviceImageBackendContent/L0_interop_test.cpp index 21b52a9fab30c..fc6cfeed3e3a4 100644 --- a/sycl/test-e2e/DeviceImageBackendContent/L0_interop_test.cpp +++ b/sycl/test-e2e/DeviceImageBackendContent/L0_interop_test.cpp @@ -49,6 +49,7 @@ int main() { [&](const sycl::device_image &img) { return img.has_kernel(iota_id, d); }); + assert(!exe_bndl.empty()); std::vector bytes; const sycl::device_image &img = *(exe_bndl.begin()); From 5d7d500bb502267a4a3d1e8a26a9f97a20abd9d4 Mon Sep 17 00:00:00 2001 From: "Bushi, Lorenc" Date: Fri, 17 Jan 2025 10:19:48 -0800 Subject: [PATCH 13/53] Define feature macro for device image backend content --- sycl/source/feature_test.hpp.in | 1 + 1 file changed, 1 insertion(+) diff --git a/sycl/source/feature_test.hpp.in b/sycl/source/feature_test.hpp.in index 78199433a9e8b..09070e74f2392 100644 --- a/sycl/source/feature_test.hpp.in +++ b/sycl/source/feature_test.hpp.in @@ -112,6 +112,7 @@ inline namespace _V1 { #define SYCL_EXT_ONEAPI_WORK_GROUP_MEMORY 1 #define SYCL_EXT_ONEAPI_WORK_GROUP_SCRATCH_MEMORY 1 #define SYCL_EXT_ONEAPI_WORK_GROUP_STATIC 1 +#define SYCL_EXT_ONEAPI_DEVICE_IMAGE_BACKEND_CONTENT 1 // In progress yet #define SYCL_EXT_ONEAPI_ATOMIC16 0 From 5050a3c754de100c62d7f38ada8ea4aa8ad5eb6c Mon Sep 17 00:00:00 2001 From: Lorenc Bushi Date: Fri, 17 Jan 2025 11:37:22 -0800 Subject: [PATCH 14/53] Improve testing logic --- sycl/test-e2e/DeviceImageBackendContent/basic_test.cpp | 4 +++- 1 file changed, 3 insertions(+), 1 deletion(-) diff --git a/sycl/test-e2e/DeviceImageBackendContent/basic_test.cpp b/sycl/test-e2e/DeviceImageBackendContent/basic_test.cpp index 69dac913418ca..821bb7d850993 100644 --- a/sycl/test-e2e/DeviceImageBackendContent/basic_test.cpp +++ b/sycl/test-e2e/DeviceImageBackendContent/basic_test.cpp @@ -28,18 +28,19 @@ int main() { sycl::kernel_id id = sycl::get_kernel_id(); auto bundle = sycl::get_kernel_bundle(ctxt, {id}); - assert(!bundle.empty()); sycl::backend backend; std::vector bytes; #ifdef __cpp_lib_span std::span bytes_view; #endif + int runs = 0; for (const auto &img : bundle) { if (img.has_kernel(id, d)) { // Check that all 3 functions of the api compile. // Furthermore, check that the backend corresponds to the backend of the // bundle Check that the view of the content is indeed equal to the // content. + ++runs; backend = img.ext_oneapi_get_backend(); assert(backend == bundle.get_backend()); bytes = img.ext_oneapi_get_backend_content(); @@ -52,6 +53,7 @@ int main() { #endif } } + assert(runs); // check that the inner loop ran at least once #ifdef TEST_API_VIOLATION // Check that the ext_oneapi_get_backend_content and the From ab765f9638b183c1ccc9abf9041960cd2505164b Mon Sep 17 00:00:00 2001 From: Lorenc Bushi Date: Fri, 17 Jan 2025 21:03:04 -0500 Subject: [PATCH 15/53] Make comments more helpful --- sycl/test-e2e/DeviceImageBackendContent/basic_test.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/sycl/test-e2e/DeviceImageBackendContent/basic_test.cpp b/sycl/test-e2e/DeviceImageBackendContent/basic_test.cpp index 821bb7d850993..aaffb7ceb8654 100644 --- a/sycl/test-e2e/DeviceImageBackendContent/basic_test.cpp +++ b/sycl/test-e2e/DeviceImageBackendContent/basic_test.cpp @@ -53,7 +53,7 @@ int main() { #endif } } - assert(runs); // check that the inner loop ran at least once + assert(runs); // check that the loop contents were ran at least once #ifdef TEST_API_VIOLATION // Check that the ext_oneapi_get_backend_content and the From e8be85939eacb0d25f2f6be83e76e9431dd19a47 Mon Sep 17 00:00:00 2001 From: Lorenc Bushi Date: Sun, 19 Jan 2025 18:47:59 -0800 Subject: [PATCH 16/53] Enhance tests --- .../DeviceImageBackendContent/basic_test.cpp | 42 ++++++++++--------- 1 file changed, 23 insertions(+), 19 deletions(-) diff --git a/sycl/test-e2e/DeviceImageBackendContent/basic_test.cpp b/sycl/test-e2e/DeviceImageBackendContent/basic_test.cpp index 821bb7d850993..2e33f9c610d49 100644 --- a/sycl/test-e2e/DeviceImageBackendContent/basic_test.cpp +++ b/sycl/test-e2e/DeviceImageBackendContent/basic_test.cpp @@ -7,6 +7,7 @@ #endif #include #include +#include class kernel; @@ -28,38 +29,41 @@ int main() { sycl::kernel_id id = sycl::get_kernel_id(); auto bundle = sycl::get_kernel_bundle(ctxt, {id}); + assert(!bundle.empty()); sycl::backend backend; std::vector bytes; #ifdef __cpp_lib_span std::span bytes_view; #endif - int runs = 0; for (const auto &img : bundle) { - if (img.has_kernel(id, d)) { - // Check that all 3 functions of the api compile. - // Furthermore, check that the backend corresponds to the backend of the - // bundle Check that the view of the content is indeed equal to the - // content. - ++runs; - backend = img.ext_oneapi_get_backend(); - assert(backend == bundle.get_backend()); - bytes = img.ext_oneapi_get_backend_content(); + // Check that all 3 functions of the api return correct types and compile. + // Furthermore, check that the backend corresponds to the backend of the + // bundle Check that the view of the content is indeed equal to the + // content. + static_assert(std::is_same_v); + static_assert(std::is_same_v); + backend = img.ext_oneapi_get_backend(); + assert(backend == bundle.get_backend()); + bytes = img.ext_oneapi_get_backend_content(); #ifdef __cpp_lib_span - bytes_view = img.ext_oneapi_get_backend_content_view(); - assert(bytes_view.size() == bytes.size()); - for (size_t i = 0; i < bytes.size(); ++i) { - assert(bytes[i] == bytes_view[i]); - } -#endif + static_assert( + std ::is_same_v); + bytes_view = img.ext_oneapi_get_backend_content_view(); + assert(bytes_view.size() == bytes.size()); + for (size_t i = 0; i < bytes.size(); ++i) { + assert(bytes[i] == bytes_view[i]); } +#endif } - assert(runs); // check that the inner loop ran at least once #ifdef TEST_API_VIOLATION // Check that the ext_oneapi_get_backend_content and the - // ext_oneapi_get_backend_content_view of the content functions are only + // ext_oneapi_get_backend_content_view of the content functions are not // available - // when the image is in the executable state. + // when the image is not in the executable state. auto input_bundle = sycl::get_kernel_bundle(ctxt, {id}); From 74c0772de4707d1ac89aee7b110577427e1d63e9 Mon Sep 17 00:00:00 2001 From: Lorenc Bushi Date: Sun, 19 Jan 2025 18:51:04 -0800 Subject: [PATCH 17/53] Remove rogue binary file --- .../DeviceImageBackendContent/basic_test | Bin 173848 -> 0 bytes 1 file changed, 0 insertions(+), 0 deletions(-) delete mode 100755 sycl/test-e2e/DeviceImageBackendContent/basic_test diff --git a/sycl/test-e2e/DeviceImageBackendContent/basic_test b/sycl/test-e2e/DeviceImageBackendContent/basic_test deleted file mode 100755 index 305782d7235e80bb39da688882185697fc4293a5..0000000000000000000000000000000000000000 GIT binary patch literal 0 HcmV?d00001 literal 173848 zcmeFa4}6?emB;<0EfAp6Ac|Egj##t^ElFESDFvs{fvL0*X)B1uBu%DmY?_2*Qd(>@ z*mi9gyBl4Uu68x6tXq}UsGuqGXNr|ZVKoYBSGzcB)f-!7R2HjN^M22{_j%@-e@SKC z&%5vEbu|2*`|sRy&pr3tbN@U~-xXVW&4h{y&;L#EuJj~Se@RrM?0-FP*GDF3SpF0B z=6f&Vzvp?c_D&`|N#i6%(xQ2IaDoOsA1iIpn+Qs|sq&y0&BJLL@;o1X4h5xL?iY2B z7U$2fK_9&;8$z!z*d(~0<;PAE&+n|r(a$4W8D>v-QIh?-H#`_p9C$j0k&T`)W z)y4zW>@ZJm)}P}zmLA`m;W@`^g`;Zb>%Y8SVl6>MD@YP<1e=HKKk{&iE<*B6a6`H zdN;cI%JoCNtH8$!SHZ)t78P|J7dP>~kyTo-<*-cQsL(emz*r z-;MmO;LlR>yqkEI-_88p!rvrVQ_ zX?OP3oxfr7)nB{wv~N#7_fNmtykz|sH%~tLq7NQ#TzkpH&EJgt>A6#X`R~8KD)u9D z#y>Yc{MFI>zW(R${I~gUpRzvQ&i@fCD>sr+xCwo>su zlS{|{=0&CBPeb^n=zQ|@((&8NwD&-n_#4X5sV&3j7p9c1_xUpPUs{Ha=w_*Q-B*VG zH_GsNItpHj&U?$y`EVI|K2)aPKU9WJT^TwTGJs0)b9R~jeXtDuX=V7C%fKx~zrIYp zyUN7hQHGyPnR?fh=`WdfOYtLyrd0gSGUH{ijGlb1jGQNy8NbuZ@bh&hlv4cspo~1< zP)4slS%%J!%e1SrjNX2~44hH}D{I`W+qLANo)7pNUtUUTy2GA}qg4LCHVt{8aHPUqb8_(8wGaIfIrKX@kWGK}^oz(jnaiyezS)G0~oEZfn)-mn9P& z8=6}il8KFyVM#%T6-%0lZfdFBxIweeO|&+(ENNVxtf^^S*U(Xy>}Y6Cb}ng*#cm{S z&gzEF=El0V&SXcTVZ)LYbt|gtV$0_VW7m^>NnKq~d0l(5qh?NC@aoyI8?RelExA{` z`T82#(YR?-!|LXZHPWUvU9FA6bzO3Ed!nwnbxm7L$+oq0#%iy*ae3SFI*Dj)5Jsu( zvegYubqy^oZHq{MBSPVNou@< zEULSK_=z}yicqp;bsdS8L_=pHP!d!;=!O%Zv3RzWJdxTiqq2skrjEp#psqC?iM$?M z*?N03)!&%t?Ank(3Z0u9TVQZ)bEjw>%Ri}UuEfr%dvndaro_hP#>CaC;Q8^%hUS)< zxy!p&Z)oo9WY{cWz?985Z&71oqO-HD<0|A3Z|h9XTboGMCOd4y!qo}cc{eOynS^=` zI<~eo(R2-Mh^dO@aVA-3&+bTUXxo@r+|}OFjA|zqE5n^K(3A5vCK{7%9ZRlf%)6Gv zmNzU}KBq3GL;JeIifb-gpXg{!wA8Kc!U`qo+FKf$Tj$VSrTMHWYF14(?V_!V8eAoxveYFm9UIZQ!=qBS-RD;XA7_Sf@CyXPtrr8zaTT{kfMBok$NY`)ddy5^M=({d^L00JKC5Bm;qX_9MU%{mgEL! z&IR}jdTmawjj?hg&bJJ+on=@g({F5p&d9-F_Pgh@M$B%tl<%VLBWNPNZgtb-L zs|6##r#yQ;qobptF)73swKm`)HR_q zbxGCvOVD{Wy}1GD>jR@;%^Ga8)$KV}_ioHhDJ2y1QuJ~KUW|6LqciX7Yc;C0()o;q z4c{f|Fq9L)Z0fj9t>u!=tJ_2wuqqo9wiQ8N_>GvGKctu_%G;!(!g&%XZwE_Ltle09 z-Nxnf>)!tM+I%&!251)gO(;^_P7twtwW(E! z6tQ7PY5v@uH!sy(vY}x;UYxA6lnE^{mvEzsH0V@cLFZI&OmsA_*}Sqd(PYQyRqD_z zk);Kmd)O}26!lO*&hOm3@yF!BT5Ut=c|c^<)!K{_$?}WI9k1iE9WLz8&gc{p72iX{c__HjO5)C*f9%Va|O^p{`C`UW< zB)B1XTH6R&3DsyeRss^3FM&pz>DsHWuA4Kf+Pik?lB=$+n>}mJ?BKCxZtyrK_c(jj zB{urzn<+SV*6eEU>YHy~bk&l&nptyZ%?SrC%g=aOIISj}Gn|jziV0o?CxI3GPoC`m zCUEv7PbVo(0Bxv(@X0)%!e6D9n&3*E>YbKX-lrru(VML0{Wwe2rpSL^l$Yw~n&{{W zT7LW|8=mN$?#lg|y&2iRiSSqZ+Rf}hy~sP;ZRkkbCi5)8TK+!&z7Mi9e7bk8i@RIH zmEHv|e7}ZI_U5{9*DIRYX*|h$17Z0|+Yh?>B;0=N9(J%Md)K-Ao!|N_^`GY5;=-FX ze5$wB4{QEYyrc{N_GN!!$90mo)rFt>?0R+vCwlL5;p7GHXNPZsx0|s1Ox+?y8sWWM zpX`6Kq4*`Ccz=HjqI%}%C;9t!5-;b3{?FfslX(BkP}2N;IEnZ7#3bI|hm-iXNQC?O z`;-!Yy&oy)_SGcbKYx^Hf4@`Wmq~>C`TLO)f1@7>`um*{FMF5%&)=VucsbYffBycZ z#4qj}lP7J za<<|Bc7@_E_9MamP<(YLeorXACKNvqik}^d9}LCM3B{*F@pD7*L!tOfLh-|)_<5oD zgQ0l;j7gF+q4@cs{3aBCStvdmieC_lKNgC=JQVM(&guWcP<&-5{>`EIs!;qDq4=qx z_$x#4kx;zc9r1rNLh(^Q608ozTh46a{80Q=q5RQM{MDiOcqo2xD84on9}C6b9E!gt z6ki{Tzcv)VE)?&dok?{M{F_4YZw$rvgyOFY#qS8kFAc^2?eT9C_?ra&CV{_6 z;BON6n*{#9D}momIrk6otw$^4sgr(mafKJ(o=#2}9gc6^U-=bn)aWHYCVKSz!~A#3 ztD-zfc%7t;WJdX&|9uHpNO)Mo-MhSxS%e zY54OJmL>E^kA^=bVOctlv}^cd5|$#}cS%^5#v|1l zez$~WNjwtK@EsDCrSM3VhC3uIOW+Yt!*7?cEPY3^e_|q#w@6r)t|P-5 zzEQ%mWF1Ls_>B^lrRvClhOd&aEKx`LHM~H=vNRp()9`Ew%hGeCN5ij|uq-)8+BJNh zgk`BYQm^4xNLaRxMrt+uVhPKVb0n(a(#iPmr)IF-IaA{=+Q7va}qj((rF2 zEKAA}Ps6{Euq-79?eF@9bab#G--vA;m0H_ zOT>|W4L>MhSsIS?Y54OJmL=gxkA^=bVOa`}v}^cd5|$<4NWF$XBw<A5^SVi!>^FAEWJi*HT+@;%aUs(s^QZlEK9ACY7I}2uq?4gA{zd~ zOv18+8mZFoZzL>Brx8!XzmTvjnMSgI)c%)nRKghze_z6~L>d{^@V6x_OQVsrhQBUh zSrUy5X!tP+%Tj2hU&9YdSe8H|eH#9}gk|Y7(xc%|Nm!OVBkdagn1p4iGg7bN4@r25 zgljeYehJIcW+bZNyCf`2nvrS^zgxnxlo^R=_znp#lW>)WJ0yI)ggwHw@kgtF5GBM` zt?1LuUs)r+SyNuMO(VnU(m$jJX<{$MzSjJ5N$=>f-*;?cyyEfrxBi?wGoWyaqfl60>hHPY0c7BHowYLF`%`lVo&NYqA19{b)~^)R zC;5-umE%E*#P2*R*Gl8P-N)iPyVLRBhi5?J5u!YE2+KueKapaN&2hqyw0aA#JYUj!AC`GwuUPa&gE9}9A>lFgfR;?{lvH?- z!ttFC*8^rPxuhl=`lM6_6DBPWckY+A-};q!@4nvQc<;_l6zJWa)Y9GkW)*pkoKrG! z$h=G0eprf_r%X5+nGCf>U&6HkGu7~$CP(o~G_%Hf)H z?>@7F8J2`0X7cE*A%$}N(+-aZV#jHRux`u))W*Z0XOAs$4gCcbm!p!rW>S+Wh9 zk4S&)?A}G%LGwP!X#5T{kNHIf5|wuLK75dwgFHLTFG+}}wrjX|d#$7%(0&kQK!9(s z`7`75Teh|j3F)orDs*93 zYk62<>jC8Va9ViSDfv@7HDhYKM8z)?)scK%--N36&3ni-cXUrh*MuV{jXW_~*sTot-CCW-d@^yW4WIr7U z=?Fbxc3INeD{jE09-b;GACZ*a?fsNNJyQBWxJvrZW%;npvdhmR4`iv-EH>xm8k(}L zLh4M}xZdraC8B_;p{mBTLX2)XY2;y%ed=LJG`&Rb*nTCCcjYN8l~&%p!yItK)SPQ$4w_5sgE32a zNX3q0ro;qbrgoD59zXxDlv?*ebCREbnmnYs2h7#Sa`x!XqXW{c-iO!2*+qR2;8>f7sWOzHtiil?j<{(_W0FWe_1dzHytCyjhmB4{ub8uNJ@_@>$X zoUpD|N$dHhxd&Q1w%?@HEKO+&{L=<)@Lig+c*+-t7H$1e1$qa)WrOD3FsJ2rnA?SO zVfgRu!#B-${#!V;Qo3`$NH3Tnoc_=aQNP$tezA{eF{J-hh);PHUXz;-B39O2StNz{ zyE9X6N>4fSl=y*bWghcP;ygGI(I1FnP~^EgYifDe5qtKCcUSCLHC{7`&d{=~5&ui` z@9nuV$}^$f_|*8`fn2DHsU#v?A26>&O{To6N3tnDOiBY4({T^h^p)gKYLMlm^si6j zC(J6j`3;jS?Xc}Jzaua8C}&@m=dJszZv9Hh_V=zlnu;CmjTyY6$wyFav(;22UvcD8 z5sdi^>xTH=8UAqYtvqtEO5%fBflQZV8ey>)PcfI0TqQA5Z))XHlO>@hJ-xy*y{52! zCJmW<5$v@p?PUZ`g^0}_me1Cs`TZsG5Rcdq4*Qc%epVjB-G#zkFEhN%@+yZs|1%`- zfFJlb=BzR3hx{`YIgV0_O17&(Qg33rx?DTE2N>klj03X*ugv;R&pZ_snwhjw#$%u_-0tJQM+wa z>H#qVIk!wKfypv#H2Zu>>wQ?KnVnt3=5;Jz7`#)p$O9UkQpXIo(DSM(Rcd8$Ce0R_ z6VNc<5FzgD9yaF-g#y=Z*^b`WHc|)7dl+FtaC#24; z`)l*ZkEJd8b)B$MG`^}Cj|Z*KiG;e2nMqfP!}y%IZ3kjm{_Ab+&Q^5SX#J3R3F)*5 zMo^^KuxR)`jaoL$kU*!lrYq3mBVRiZJ1RuBb|0PCeXf*9bsaSyl&l9lNnkW<)}u!b zsX^Mhzq-hO-FnoQPu3rURz6HODj(QT`5-58Vr3dCCm-1vDwNNiN7E9kU7)iK(#X}+ z`>>sN7-WBzY$*+!Z?lFp8?Gm!_i@#k+vHiMz0^va>F$gf&4srZ zZ&dF$tY%%R9a=RXq_sqMkGT}a$F*-#xmx>T{l4>Pp9n56 z`La-p^*kPXRyfPRzIYSn)f0sM2jm$&>51QYV3#em!+cSbY_`)TO15r_|FFW``YX-6 z!#uZADqspZ1&w4wvp3cU?X)>gy!qanWg`42f=G4sN!hz%A7whXvw7MK{nE|kY4b2S z%`SEeY^U~{-SRBlY&}p*>3%wQkNF+BkeoCsC%M7i*!^1NW=T!q4S3Ivj^^3-U~Ds5 zwk{x7S^Hj}(?DJ9QLpxOn)p5D-&91F8LAvL^Wa3gIc-jtmEG3;^(7d~vcI8>twn12 z{n^eUb61@GgA}VQReo%ZxOGgh$g*D=rd3~Ldq?=$V=m5Bd$-Wuy1%_dwVVBFv&cC& zf6*R1$yB${M<1A0v8uib1$!KtAH~5Ygy&=UWmT$cZcTsKL2k|2+ z0m2869fbmSZ7$Qa>}TcoSI%dno{YI+g=G(Mn)qPtaMTe=SH*izVZ>CLKUrZ8I0$h zGOhG(ZxUKATfby;v*AmAAqpSX7gQN9(Ww-JcN>Ac^$bd=wk?C444YVwn>OWFmng<` zpj-z$+1PhVA0AJ|`g4}eNUxUCk4T=riSE8e^6PR)nK0%mjlikTdGw0|egpk5my<25 zr>CFcd7(YoJbN$xRx04Dn183QMoxR+C2MAdGUv%FSO0~fEIxm}BHBSL63Dj?q!-xF z96jd`>1gojg~pS4A?c1E>g+3$$8{z657)O|(#3N`BJ@mP4?aFQ-yUG_L-cL`2KLXI zdAwyMW!6gASvX4+gOoR){O`&m+`dxsES6VDU#)t^w;uD7r=cV?>Xgx1w>U2$-w(@X zwzF2H`ZL8qDNe8U$A?{C(*c?IRx|PT4Rk^#o|^R4YK;3RPiXWKsUdHXeXS~Q&99Mi zudGeTv`y?piUkad>=|mjb#n+EKTWcwWI}%2oJ(gPIaLeD`q(YLtF?3T^gkpInG2XN z=15Uz8f3oMdRk6Dth})=R43!LxV=Iz>#rru?Jp)zYOM8qPCu@byk+DY_V1oL$Iq7= zjQ)JK^{AD%^#u-!?ZXCU&7)FA7gaQH!ZaH&7x17hMJ}4^)}5v|cGx^wskw~#$HQWP z{rL>~;$|TJE^sAHUB;v%T}e1a7{9J}pLFh`-n6+^iIBTs@#WLYTOYMc&x*!MYkOB> zFJ*aZU3!qX*bdk<=6%wDeVGYF^hr%1JNVM^c9z+v4YdCLPFa0&T=L3SYIAzv6Cx}7 zIOEv0G_|aM?{P2g)Q9o!FOGfe3KJAyF!FKD^?i9#c*^BpHTg4ukZizj9ZNxMG zCAr!j8#ad`i#U`(tIjaJe;PJ<5Y7DZ-HU{6Kp9rCN*YXUSupv&|eb zpN3FQql)e8m$^z&PCsQ}KEXICPLFDjTKjp>tW{=5J}E=TF_GoGmtSe zAFOVEmGY4Nk>}&E(ZmcdS3ZMNUnJ z%`eF*eP*W9qL9B=C5=#vAUY&${v?=Za@y+Klkj|aVFHRp_O#Yn7ih`?Q7nvLH~e;G z%*&(RIuG~8o-|c!wB)e42~Eq_ zkBGL%jtA;K1>)F~Yb(L5*|Ap35h22T%nYs9N@Sh(u+$$cqMl_P7H<9GQvPIgSX(#I>G5_Kw)f3ExwL}7~3t$h3$!1cN4tGSN+{uX+v z{yzQtqoX@wzmr3W-&wjcRKJQ{Dc!#HCF>V`r!!FxmO44rqdU-uV^k4n#1m3pPSSUF zKTGPc`98g+svs(oS~FVURrhXl4Rt^LG+C^9mW5)c4BnA5Q517I6^y|Vg{1tdL>&Hn? zjg$WA66t=2;2rcXJG@x^4jhD(tbr@V;>4LAgH>P#oGM!c!TK@KKix$*MXbN- z;0eiZon%)}*^W=yoN@BoU@6fPqIK+{mXTksKgGjw`fb{R()sIIX)kk#zaBBSe~;z$ zi}_EWAG!93-ZYpr1;2^!`@NT>+V~c7s zY$2KX2LB|e(CI_4f80Z+9{jgWNJe&p7})&>V!H{A_AESK?m&6T*P~$lsn~924c(GW z#RgLz8-l52_fvXsVx=5U$b}49m-sn#cylJ0iXCD`W_Ha^*~j_C)2Z0gi{V!~Fy(pf_u{0H}@^`6A?7v_< zod<9IpA3ibSuc4eijvW(P0aa4wl~VJm)|-w*Il{!t>pO8`N8Q61jG7Nm>*4GOL5=N zCYdL`qw`N*e#YAiUB9b7-Yy3ldG@4UCJvoQ)Q_dnxSNlED=H73IYoF-7efXw4%s}$ zHk~Y%A#48WDGg+p9&p!iPX;l|dMo8?6qZ=?Gh)bpko*VD@BQR|l;l*G92N~z_Z~AU z+1O4tKYLo;%g_Qq@|ag)9#XOWx}p##{5?tlZ|r2t!k~4+pM`I;91>L7ua*UWh;tF= zaQBPDJy&Zm|3%BlAw2$u%12~x@yM@yFG&>DJg8ZYoEr~G#{WvASq2$1 zQs)8lJGRG5`BT3AD!#w)uiA3fh~fJS14Nd&zwm&cfA2tj6e);!sA;ejm(zDY z9lNK6$K3Bw8%e!Pp68INjCt2E@>Xx4jQtMTS1FM-uavHGt?d!9>H5LvheOk? z`!K5Pvgg?gm+sGRI!?r991sc|aS zFB}W+XW0I7`OPWP`|iHT6T(6H{u}$g$U9G>@ZO&9Mj;|3{re)H(2#7G(EFPuQ>ptR zJtUmaeG$9AD{hawERs`-u>1r0`sX=e`N;m5>+3+AzP!Wp^*bg>Lu4xDPvw%UESJ9) z`Dd~J=#^xZ64(fW))WmGOu3k%0n!IYWhF26r_S z{Ppb^{`0L;Q7}fQ{CBvobM`H)zhVE`=?`;xF@GWdnfH@|{#QK>`^yu}C*c8w6m#~B zWe1^(neTb-f74$UXY&O7<##aG6zgkFpHv=dmoqQO-&!9ud0ctF0Do(udFzw@oD!xV z^he3_AoN$eH0S?~n;&fr)=T03Hb+XzFGuf*lIhj{crAB+Q8^r(U-aVR zY@M{R&M)qgY=8api?47j;@bnO|JY?z{6x+#zVv%xiTS1MeZ-vnQj)i5zI5rnzp$yw zOP(zzFP^%l>Opp~m>_;9(}2#>{JGv2Y2p0v7meSLeLf-nPx&A4FSkezaFObTwqMG>=YLrK+Jf?h^R;!D zd(atw|D{krnPU;Z?jQBWjLcr*da;YRg)%hF>E_S z^BS-VY{pc<06y2xJm|B+a+V!Fh!{%QILdl)FiZ=M`ZaHjwL@E4-P!QF* z|8bcfg8k9b`t{9Ipoo5Hi62YAZsoA>c+5;dNYoza+Kt>+3+$1l=lntQGC3oJjxYdMx#EqwpT z*N^em-&$p~a>^s#)?=^lda-mXyCkg9(iPf#wbK{Yr>`1fy#xbg}zb&GY7c zA{0HQlAQr`W(r5oPBZOm8dmr@VdT&EO2?c|E3K`|jpwkv{iD=UG=GQXnKiege4+#K z!K}Zu)uVAczp|RCdHz9R%)D)%s+TbwTS;H*!eWvJ{!=^U$2x}-?%dBbAQRMSuyJQ?O_Kl+}NV_ z`1Q+L^W(v~w^(0F*t038hwQ(;KPqcGap0DL@*)A9YcFG99vIZTHY4uh){%<%sYAZT zYp%;vgdu+KBmTD`;}81(DvSB4Dw1q)&X=jp#3Zbjh`KXg zm!p<`#HTi7=2~h)C8^zGsb%P;XuiPm`^(y^ta+Uhqk((<&tE!)v)!B%ptE;@e)?a< zEYQ!w{?&bPyMZ>%Di2+oCvQ9cRNm;dJ(-ZX#YBZQIjWXd(B5DhosrwO&Oa7HAqCIv zpZT=JuPvO9t=*A+q<`rYJpaopI=LOD#`@S1w3igQR?R32R}k3ye0>*v8{1zGwNLcw zuXum*3GOZBBw)u&*xsc?W<~bS^6&TmU*b<(Sc0Cf4+lsqV~+~@JE(s=f9D0IJASCM zSRU(3@Lwb^CN`bB-438Edm4X|v(JV4EppVw9tS=@u=LT(@ca{=UyQlw&kSTeVD->4 z-iMR06`pFIP9AyV<9}CP;r5?iLO#Xv%;|HmKebX0Z?b-Wm2AI^$soVjR^-K%y*#(- zk3TA<&&BoWz8f|{$|Fl`0+n;mCzyJ-@f@rtto+m_*uENb|5Vx=wja7JARQX`kHz%H z8vl;o3)06gOOs0I<5PNp$hYS?`B=-v{B(l%T}$=%c=|A2da=Iz-__p}lUF@9d#JaI zqvmBpaxl&*A|P89Tf3jD=subMdfi=*M(Hf`@-OM!JwUJRv2MWYSUjSm9GYiD|J~KJ zy}$Vu=ZqhIP)`s!wTwUFf7Chtp#RUunJ4#Z(LMgs<{*da=31#ijzXW&RE{1RgiONPRkoSbOT`&yw_xAc7tHYSPc|tKgavFZi4?AP;(9 zNnfI#zybZE-cuqfH2aY&+|z>*3GUtQg68TkN*l!)5Z|=C|7hnh=5mJrKL`nR&SXO% zXMdeO-o-~K^sVc_VJI;_TzJ%PYp#5rzF$Ti1@oz+FZVFUKabS;Kkz;5dp=%hyzJ*cc`HZq7oMIWUwpo$?^J z&Ey)WyafE1zchb~P1!-p4pREe9!=r?=1%)Aru`m{`J9r{Z{wI4&%N8{!ivn(kbeoQ zE-toIxfp$y*<7qw<@Li4sRtx8>v$2|(o}pZH={FCz8H@zO|9(L_eAxh8B2R>qkJ-_ zZ_1g`YEsgKXAtJ~(I_QW@`;98-hPe^@jB_!-e@gvK=#JtdY@p?psf4MefKL%mUHtf zW!b(b?jJcJo~`?)KF)Ur;=Oj^pl8NAkBJlDf2PBItwSRpr)9FD)(@+Aw4&>W@s$0V z#b=-|0#I4aC&|3m?U87De5ZZc=^mSHr+uqhuMy zvdhxDJqvdd^RPUa@5+Xr<$*N)2*+8{0V(|;2RFH`l02@o_f8Y0ezt}Hllk$hD zVKC}+z&+Ki2V~01_zV56&q;%E+n2KNvco0MKZ^t-B5aS|Aem)!>gM{JpeWazmQ2;V z7CPpvJq_EJ^`bHb_NCC@wD)-fpA+&Kv)q=$?iG(;vDa!Z^d)b9Ra0c&cwb$*igq3| zb2*#GTu;%!wKg!}$ZPLpwj)c|spOp~efVNxq)^u>#7vSH{>ODI+l0Wr*5>3Ra+cS# zL-q?-P+zZecdT!g?JPYzm;2TDc$cD~D4(2@FO-`@_BwLm?VeXllH9Ex`n*O={(Qkm(+0NWcZD8 z*drds7INhjD(KIkeelhGneyd!?Kolgr)OKk?3^J#>{7nG8$j0dc>1sl~KR z7G2Ol+SJ*A%o2Od$8srs<^vkwo%=)fA!~l?$DBro)0C{KkcT_>OI6|=l+zC>SKMF9 zkgQ#F?9uIZU%^hIdchzgmkaVopWl8R-=g-%bau2eOM^7YpURhqQUIgmqTx8bk}yW)TrOD?MvIXw$Ohms&`^p8|y{^+WW{njw2f zsNYQfmguLroJf=TF4Z+;HrYaZM6G3JT!!{yQcAx+A^+gKzTd`px1Z{gGj;Ryucw0Ert?Dq!#i<#c`YFJ)szgc&r7FEig_XR_IB6CvSen`Zg z(P18av7$mzl?p0UxSyr;^6sxx2L=88Avq6}xqrP9(LmPxy)<4Jtk&siI8Oxs!YDwd z88h^$vF3ltRbb!!_en(`g7L0;aS-7UCmkx|UcT4E_mOUrj#VbY`ka+sdL`X7(|yh7DUFISb&N8IBAeT;_7tB#0=h;Hg0?Z1&<{mMpdhJCy1n33;EA>bjo z-@czLx_==0toMa=pFqA$g0LU(2cAr?GAzvgk6YECb28>VPs$u^s%iI;YOTV(Q)&&C z&dDE_L>bId`fR?)BNN%5p9mYT(0*&4KAL-^y(te?WK3qwdnBxVQu;o5sqxS#Wy*|C zXP*k@uH1SlNXJOKgW0_C>e55wsYw6EIWH>6@77BB_q%GouWFU&bK3DM_4{*MaIY)R zUKh}3K#jG3T}pmq=g;e}z<$)(FB#bp%zJOr0j0)kN&eO2PAzdZL~g4sp%qvU8JYg( zihZ`12K1RqO5f*Ie`3u2I+@I}$dP%EF|gHFaMd$vKQrFTJ7I5vpPMjf!pjZ5B?}CG zFrGyo`SdzJ<_63QqG$V0`1Gg{85Y|j0q1C zC07At8OUp1aHdYLM#!Qfa!l-LKhU+H6gOOag-xO`BPMisN?|DZ5DN zH`{DVAD@GT-DPPum;3dzdP`FimiA7XzBF}J<5>%^H(76TmbAQKv8B#hXKCwG%Kj_!> ztlr0%YrmJx=g}Z$FH~iZd?!b?&_8$fp1WA0m{T4(l^5LZ=EbV8dlB~cq&K;%m7g4} zry{CL;H zw z*{(fTT268TKY{jqpNj$c?b&t0?J2t-YsbIy-|hF8L|JFyFp7kw<#O6n(J5E&MIr{t z)(R(fpo=^t>I#h#Z}p(;)mWp3J(TjI&Hhcj&s!4^%vVwE|2w;nQQ^BbI>R0PkkTK9 z%+RAOS3Gn5hlwdR+$Vhp!~NJlk7>AnD+5KH&Cn{gl^XI4w~XK7^@@K#tAM}$e=6kf zGs1Bh{>~elzvZR)`vx`S@ptQA!k<^f-x;O(+aU6LH>73K!5=R*@qNLHo<0r9Z&o3H z50mSJMg&nfmVd zS(*CYR9xSmHA=tQHFkXw+GlN+49Q2M2cyy+j~RK?Tp$m$L#6|+#lFJ!0kb4&GbPQ< zRy{oX_Al`FOWJeapLUiboX7bqOtE4{kYefgdD}_fWBX%3vOXUBooD(^#RPns|M0yY z)r38%*h4b+HL{-gIPpig7M>92y5@mJui{HQSk`#&)dbc>OG_U_xq;WqXi}}@$nur@~=~AQ!-{JkG6kg zdc64IaAi6FLhVeoI(38C;^|TH`pXZhMX^8DC{(%=7HH1VmW}^fo_MbJ@|7S~l z{mIW(RS zo%s0vF>CSn)7&_fgLoOE=0}#&u$_v+{*BJxLXRPuqlb7<*_8|R8U_LKf0chjuIOu?Ua2})oGDs#@#oT zPu9hZdAoK3E#LQ}3HCgRGT! zua7T$pWjcp@eocgz5ccOVaKcJ#IySTJ&i>}A2)74E~HQOIsH&R2f4ybmksFa5-ATs z^h4#K>3&uR8$PxM->D4|+bH)|25~dLj+xvOPi;aU)DwSF96UKOknfSXTobv;ES_*B zG}7jdFv+xAuZTwo59FGt8TG6oSf9|pHR-(lBH@3i-2JzL^Sqb$!IJxew7dT%GA?r- z`Qd+bZba$xyeHqUMfdo|-rj$c8f4$_8BKNPd8tW5ba;JO8v}o~Dkp!L=a`@Le5=g( z4EMj@zbx7xiRRHmzfGrfl`LL7;SBMj>Jnk3qrf5aCYwV)>nBEj*^qP*Uc8r#XlrV1 zhP0AM!=*8|5~aJOZUgmz`2qLha^}sqSGIjJf4^0T6!Du&|8&XpNUlA~Eko!3NB2j? zp6Q~pKwm_j=h#jux(~jb`%I0|Mqk|lb+KA6|BKn>e- z`S&4wn{Is`)Km4~pgCO@v$|h(P|iE1#2@`0s}uQ|+V|=(9M=6b`AF83tsf;GT^=@b zbqy>RdSnDJ-j?Sg7#ZSFVSld0`0(Da3>Fz7kJ>yWxKGRR53{ZJ%gIL!_!?AdRvARbSf3ruQZeG@e<+6`1+Mu7py`UwbJ;!`NhgVJb!PsC!sP? z2HUy0{;uxrKAP%2YG&e^vg^cEU_09R#?&!!Q=Ig`gDD({g8j9yeHo{}>6ic4-}ltB z`FwrJ>FcLOzwmXJ2YOa4FV#;G(IWa!&X05bmoMIE&L5HYu7dOCaq|27JKBFbzhb}g z&kOHYHZjj9c3xN=jEBIdvh#-T2b$m7_KL4kq#woOF_K5m_l1ME?Xvs3nj+rZD9LNy!Uy)jBBXkCBD&$;~eeN1G`U6jigRV7~o=Y(#@%&*LvWhy68d{#SK_wgmccqLbbW#2)*&!%~G zMv`SC{*}#7>ZIJb3k_#ykNolHud!r~bZbI8Kk69Ewa@?FqFVw<7u#jf9$iw$AN;vL zh5pGhg?lSJ%dAL$>#R6z&a_80X~y<3R{d!+V5dUKJVgB&Gx`VVCfNy*$&^36Sj*YV z`q&WN_P7`;v%u~d4W;4p=7;rp&K@r#0+3LYF8JdWc@KOR$piboN{SJ>Xz< zgb;5M)zqdTr+ojwn%-QO2K!ICx4@o6M;Ud<)$~W?e4ax`SSq1jU|Gu1c9`fvGmFkz(E7}wKe(r?$;HRBi!S=$i zxsF~vf}x;WbUmuWL)?eedsuk-_TJJ9&By+HR~2U8_m`c#Ny;xTZT?QkiFkBdBB$a(|3>oo6H=GQj%~!K zsE35q@I)Scx#MoYGcOe(NH**G1|!(&%Q)kS@j(BF{4vg)y)Jp3AC=&nb7ISQl3f|= zpqXo=^K~v1DeK@3`!{}n`TZ+>>-@Q~^jj&Z^h5<|Bx(h1ow96qD1x zPyKuH?xU(i{&+c2etdc_IKMEDzU1FG+s7+nmaAM7^7%V~@vQHYUdk28aQl3JiGLY{ zJ16q*pE!$=+tH@u>^~^cXEwh}I&Wp43^^1@2JPUYLFQr4eB%x^Uh*%l+kYel%^oyy z>tyxB?LUrahRAxvqu4I9fmgz$ylJuzL*_cG+Br*)?!1uM88$~C57Y=q~7 zS|rBE7U2yB!qa~nl8HKQ8A>CtA+z3Ze8#j(hl6EjjH)%3ko141J+@}uf zSEE#7eW?(K{&`wez?A5S+aG5vom_2*N0-d<$LCT+@$;$WAdWrfA{e|(u zdGdC;(P&3vN&9d2PqXST>Up6~Jj0A&8Pgeaj~(E{2vz3QgW!O9w_P(0$br$x_L&Ua zj|wvLHY)#TMYQinT$N`PAPXdQZ3e*%R2b4?N`}lOmH_d&ozohLa|;GpFX<>({d4+^ zjnAKsa{h(Y2Wzj4|NSR%sX;dZns;;?4%v;?}2eP#&dPfKB?|YH)UzA z`uXNk+a*NOnWtGrs4iMBxbdvzP2v9IBq;^Gzj*&qJqY-a$xhB0b-z^}R$Gd)z5`vp)ahy^mq#G02=G zKN)jqHP%k|k9u`?fbL;K=}El-~Zb23K%OU+SU%13@hGkLl0n*?%Vk?)oF`SBkdjSh8&&|KERE{{f$`T>k|2cbxmT(mvfhgohV8 zA3qVuC0GBL`09UNojNBiA;yk5F@-25mBGH?5{pSSRSN1nclq&0tCz)ODqVtW;g zmwbQBou5O{?WYHG&E5BmqfaNW-V4ni1^oXH`s*A!y%dc10{Wq`?EIYqf2c6O-(Q~j za+}`=x&B3-CGE3l{5kS!mWnNxKYU&so*(4jfRQQkb~{RB0v%+uh~JQAn&SIhG)|`e zm1**k#3Jb(lk5ZkO)z=PQ4FmbZUaAHw$GDXU9m>}{>; zO)olK%76Qwwnq=M57Hj%zjcTKSB4)q*GQ+y@eKakFM`bKzfF|vt^Yj2+)!s`LG&Yl4>?0)GkFg!H$1O#g>d8EGnQr*r>c1&P%T z+2c_Cs{2L$&tWO`*YJPzK7+VCSE-bcuKGVJ#xnlTdf^_n%KATr_R!x?6Q4``o%cvp zWBNNHj2@a;cE0&*`8Q{ZV>!xerj4}9T={nG7;$Sn^Eqj!b!&nZxSij{Ui!YvyXf^G zxnRF8$U{%j@>2fIQQZv;W{9%wHUId3%D>6!r#*~2hAP{>Ju&;q`+rV>5lhx=lz{j| z+|v^UDOisb`BO;ptGoq z)#&@p#13fSOE$}(`G`HzVs=9I?~{ORh-1`$N)Bczxm;)WyhlOEAaNW-`^zgHwpYr0{<5zP~qWI^ZWCfsN&gcEnebqg+oo`KUZFF z)8Fi{%i-M){d_(A`%3Z~h+7Y9^@4NI0b*TneMEx)YJJ}NKIOi}OzZn)>w zUzNYiTP>gS-Q|X21_FkCHj0_>RyTAuH`XN+oyl2^?d@J(6TDZi?hO8aebn-DRRXX5tnIgB3b8RNne~;)6mtDj5Ks^Zf&e{t&TJ{bgf;N zj3hQSCfbwDZLOUPywgtgmb5k{HbvUHl99GGk&cGewFxiC5ovCXymNKenl*_IFXF9f zXz5IN7q_%EHncP~cdnm(am(ti&UMWfx3p~BaPjJ{=9Z?5S2wp_JZsj)&8>|sT}_FL zrPD|e0<+f9LJ2QeU?UdNpeq+FaP((f5P9p0$Q9I>yeP7`c|)SLQ&=RX=Ay`=7O6DZ z)*(484XXtU7BsXpuWd~<)iwB8l1wlz21>LoTNIyOfdp_uIGYJ^gx zb6s0kOH-t=BhipdL>gKn%^TWVyyY->!|Jyu8k4gm9C__)Bj=r9+4Ek2XlyGkma6Ky zlFcohDkWtz(v)ayNp5aW%z&Q-3na@0k=N!oo?dqhFClq`El2jqYmug)FD)Gjw|Z^u zi4GAXD)QH_Bh;(6dJT<@iOx>M+}P0C+Lny0PDEC>CD%oec2iwjYs+RN-V|wVt7~p; zPI`-yjK20{B-s|}OrYiHZgV@m*U+eX8d(FOk^IYXXWZ?TeC*m{nxC>k#sd}M4{U{7A)vUbS64BCOkh}x2%pH zk)cCJkd$;LiBW>=>NYkfZeOsVv8{cxcYf!BNL^ifTW4g(?dzHw*9o)d-_dnJz#Cxve-T3NEHZ53SYmKxtbchg>*d+|wCM}Id>Th6`wX6N^bjQYeUfP9gui2QJ z3ToyCma49WUXR7Rb~k?vyZqy1_sM(nzZcc#{&tILhZ%=ex7qlp!*+*#4*hwn)~Dkz z$E=&TW?!=WEmtpHme|1X49r@*t)(f^vE=&YHFY(Y7Fo0<&9NB&SN{BYXP_MYn#)Si zw`GMuzbluSYul4{{0}?y1}?c2bKRnOyO$WL^a-ed!DB!%EU}^xn;*ke*TwYyzfUf?B|zD zK=OOk(X$~hsL#(Ab^N3qW*pWRgnjxwF22guSM9Lh#rxq}7w^{_@aac5oZwnC;P@PN z`0H;)3v$|Ey_3t(Gd6v|)idnSuh);Sa4q)JOA-x5PMq)>PA4n8GrW_%SGf2y@wr9r zf+sl?`Ahh8@N{COA8%++v|hb*7N*pr{6zKCMb1;8CHY^gZ_L?2+#zZ*oB$CsFc662~a3YidX~)G_lDmiRLY z8G3NBEPUgyr}-Dgk~M(n1LfZGS9Y z=*ql=xYzS{7Mut#5#ohM(E*`3jgX8pzx=txN$z*l>E}4@+{>#OGYo9HShP;>~@mEm0(0Va{5-0Se zo!3BH=t#c`ElIzUI0;LeeZKyLMCSA7^CJ2yzc4S-7d|f{cj09!co}sIotN_`ZIFDD z?$hK^fBtx0Cj#YHan8ickHe6BX)OAt(a|VOMc+I+TGilr_23M!2b@a$0C*ADe$(ja zmKC0tzIk+Xh;Y?gMn|Wu^t{?zM@Q3d@x0-;aZ~};*Nu+W)1mF){98P)AM7DK4EBSy z_2dV8zzjG5q99%djDS@Ql&9RlYWR|HBjvzqP!w($te-=;iTd73eF^9j?gQ)jhd%q( zP;M@Kf;$LTuN@s70QnaLZ-T6|rK902c!ZU&sT5IT~61^l%_AFKtBgYDq7D=7yq1P8!X zU>e*4W(BXMevXWHgVo?e;4rvF&gMCkJw!MH9tUf|X>p>#b?8AqcrSY4aYbMU;ac!X z!r6||(LIE#xRhhSYVbMYBi!{^jNbHvl`kP5*h09y3%w)U55~bXcsKFUjqoG+LHPzt z8r%V9z&)UMJNgJlKm$g>%9kQHa0b{19;SR6j1$g)liy5zo1jOyADm0r+YDX85pWL} z1&6?T(11N)B?BxC&HyvuD#}HSPxEm z8TrBKU_ZDR90r@fNH^uddhlMb2iy-1fQP{hI0++=1*d~m@1#6f4cwEXz#i~kupitHroqFYcPI2PB-P+_uohel)`Lx854a8N2cLuPFnBLvj~A2nOFDQM z8~`W10=nRIFbggQBU{m%o8Sv6MfRr-NB=G1$jjly_5a8f+q*0k?rwz0d)x!G}n12lo^12M;Wgeit^xeFbysSGhh?wy_fP}1iTlFg8RXG z@G#f|PQICTfs-QS-w7SC3S10UgH2#9xD9Lv?*;q7{a_kA3}(PducqAJQyz?fi@|nq zF7z|tDzKWfv~6HNcpvEXL0|HN+rWD8UT^^14-SKe!7MoGJotFu=;+gw>j9?|?gtlx zX|M^*fZITiuZG+UR)hP&dhjsV2Tpno^ug(123!ng!6wk-PI`~TgZ*ISZssd63TD7| z(3?(uU;e102IOFd!R$O78J*<4~&2VU=$n%>%lD816J*V4j2K`U@f>5I(=Y<@G$7T zj`sf(bifE01#7{2upaCId%z^+2Ecy8888icAA$}T0lgW>3ygrZU=(Zz>%l&-2OI$V z!C^2BX21+s^;tQAA^juRX)x0dJ@GI81$tmRmaydSEq}0qa5UV)z7iNj@+_ zI15I>>Q6%ttOt9*KCmB5gK01eW^5N zeFi#&he5A~{GWvm7y+YT6s!kp!5**&><80e8q9zhuxfzvv&j!ez#cFPX25#T`yA!L zTCgAN2h(60%zzotn?t_OQyz?fQ7{Ttts))lC!7W66OMj?@`QW94CviY`MKl+BcS&H z`N1kM4c3DhZ~*i!As-k4y)RM^*ay~w17Ht04EBTdU!pwN17<*P5PhTme$cxWKKA0D zfElnB?0Jy*dGHBF!1^yk9~=Pd!C|mR(!n%X^$>m!<+5Ogu=nujXqIq2=v|6@z$lmj z(_j|NfZik6J1_!Hn-3r0La+~91rCE-K<`n;H&_iG0_(x!-~c%7GU$N|!AP2NU=-X9 z_JI$98SoJ3?V}tR0jDjXAHjv7_Za2CDsbJ~s24nRIq_d1eIfb5L&Oh&s|XK+Tfi*1 zThjMa4;Tdxf%V{Vum_xW1$=-D!2xg;I1Fw9v*2#f`zrHG9qj{mUkP8}Lts022b?8Vw@DSJz9tZotX;;w> za3L6Z0y)zJSB%7amG6*&A2%7flFDKGKhLts02 z2&^5ZJlGFTTTFg%A(#PIf!?<$4@SY=U_JN{*asd0`@!SjFt~#DM!pT57<9pfU_H1B z>;boc{ornJ0DK6HJOv%F7Ca92fYYuaAGi>VK23SB58MI{fV)BOJCp~j!9!p>cpU5l zr(KJ@z=dEITm@EtmvEeVz*S&7crX0ceh>d18~~pa{tr@*EN`b z@CPmg>%moEAGig~g1f=0A5tEy1rLEe;Bl}YoVE}{z2kXH@U=Mg4><6daNPci3=#5YwjDlOh zesDK906qi`gNH!x7nBF9z-i0j16&B!f~&xO@G$L-9)S+w9&k4}06qkIN1+2&fycoJ zIBf;x!G)mwUSh$kxZ_5zVpCPcIg?MR?5m&*@zeOLzj}1^wUTG?DOD0z&EFQ{PLeo& z)A_@6c*m|_PL#+qs;)U>>Kmt=etTt)_l8$o{`xu7&m&f8waCM*qoc18Ty#d&-%Yss zMJKA-VDs7HI~df0Vz- z809C^?!J0_g}n0D1?4XyJ=##7|D~kYl3tF#7SgL%mrsB1IO+Gf^jC)YyE)YUampVZ zCtbenG8TUqx%8Rg`q$;^UrKuW80~8z-D?~jy*OO{?YZ*rmGY#Qliz*gq(3=M`q6RH zCr`BX=k^?b7m?lvf3FIs-;zsTN_v*`*M-wp=h9n9_nJmW&j_b4vFTHZd#@{> zmrlIs?jF*nQ}#9CvZ1jvMEYFP#r+J%R!6RGL;4ER%gMHK(%9)UTza{-$6fkG;riF+ z>R(5C6ZOlt2?G8mUR@zMqpP z%cWOhyPhWfv~c-YuKWzrrF^;i<1W3NeAbbE|Jw59cewN`!}V{@)xU@IrzwACIDKg@ zeaNN1ES!F0F5S5Fa{5tuvgbVq{R_h7n{wr6kUnW0=K$gK<+=2@OV86?r`yuM>m2=Z zwrht=*ZIhw8*a$a-$VLz>Mz%SL!{3o{f;ny6RZMB{f6|rNq=oPy&*@x@)Xazm-KS{ z&mdjOmm7a^SH4{OI?|^%moLAA^tq&$)2}@)y&U}^(ie|G-;ln7^ts{ow}slzB%BG$O6$NN>mbo`UPVtsk}?{MiCgv*EQ&mPj7#wb4|^^;ysKMd*jlP=pt z!Ti!5(*M)ZzZUkr!|4<10{&-^o+Q1TeB&;?tofu! zKMjpT-;9%9IoY8ty@>`LUZ{>>^f5`Xs z;q*intw`|CsP zC;jkP?RVv`3zwg8jS|rDPx{@h<^7i(Ehq`LuBNg{ zw~@AqvQLG1oA8Z5ckU&9etUV}Wk2aRlU~kt9d_yE#`Ywf>UESar(4rW-$i=4vAfu% zm-DTfNFStpIoWL^eTekS!|mUg(~WycpL$z)-QMrg%Z>fRq|X?m{3JSE(#z@YbeEpj zccHmwG3jIR*ECN0wsF$$b?FS@+}syhTkI#ja*XyJCVf8X7lzA+)`pYL!ha?G|6%W2 z;Nz~U{C`OkM64LJ@{lT{RxAo(X7WaqX`5-AHngEhYmvfanwd01^B5+Rv;o8@m4{Z0 zh=?poRw{zv8Wnjct68jwf^K{)vMajz@8W}9+^x@LSKR!+=bm$azx#X5q-nA2=i_fz zrum-pJNJF=x#ymHe_^)YGSK@<$Zst}&y=C>DMQ~U^w9EyzO_X8o&x<2(3{Qj4?F#< z3X85sKo9Gm%RqkubQ-I+?euo%{Z^rey`KU7mv}!c|9j;9Fx|Hg^rH)NbG2srIVb)7 zpq~hO*!xeFp;w{fmwJDh(8KcI3OeO4Ed5Lw`kpfMeP!tTg&vl_r$B!e^&eJ$RmY?K zk^ivwmw|qNiSo4yJxu>*gdQe;59oJ5{tmN!D|zywKDQ6_XF(4ezxzQydU$RQ&Fu6~ z!*}qM(8JoN>I8m2jJ{0h_;KVLzLi#?)0U?_K6X3(GXwgw_2Koq2lNvg!ppZ0^yQ$3 z`8M{0ehuhhYp16`zX|lP@mYll`i>IxWuV_*g5C;xU(?*&1!nmx*)dP;mjQkE$lTmI z7o8^%HUCrkpr2#XFLml~pU4kupZ)TF82u@s*O~GM9r;z5*q0;!H=6Xy4=D!9-!jlC zf3Gs>o1OPtg?^DquRPy=KLdK-Xn6nL1N!}-hvk2t(8K26{h)8%6~1I^^*bpmv}$SKG_5M?y#tZBvDU{Sx$5X8w(}#WK)OyfoZCY885zZ!rVJ{j3eLJcWV`T z*!azW{s`U=(?5Gae**NdwaY%xe+hb6{`Z6aYzg^KfqwMm;pMAB$D;S6X8CqH^}7so zdOu8mwSwLZdaL>VyPWqkp!b0uR=z!;?*=_gf9wPOe$d0_xBWs7tN*7!zvG>Ab76Z6 zRj)z&fF9O9%Rqku^f3L=Ds-q>r+*lG0~w)*^{+i;(%&cVH<Nm7%u^9mzQ5GxkL@p!Y???Snn^ek{EH_krI0uDLn5 zFdg~qT2lUl{h;p#J*=*-ZM8h=%VuiQ?)0S(4XX}wT}F%W$0g^Uue<|-{&&W zm!p4&^`BPIcjNug@eg_*-VZIm(BEpNZ>*p8fqo6%5A%KQ2mM5(A2fbwPI?M-D!+>k zoa)YcgJnI3x{J`i^}h23PTfR=9#*!5(8KDc8}!qWcB3i(5=Z_x=+xhDT{T9%%{p8gD8l(L|-&SJon*_Zd^tYJr8@}Ob&>sXntbDUV4_kX# zAUsy${fN-vLUQtN>}w>1?iyQ$Z@3%uQu!M%L!T@|pB8$UZ8;12Gid+N{*Oi1-cscU zo$?$3`;3>NPYS)(tbfDzISu-@67pxu(CJ)IsrTu;5Ya=+4|=Kc zb(f)!m!VGzJuLlc&`*bL8KysG>3z_{>feImDn*Z!p(o1FyM-Q>|8daIM*hOucM^2c zzgWUJYr~h-#7H*ETvy!d{{6lhz>R3)UOeD5dnQ~iT4vi$29A_Z|s|Q zgMJI%KgOgR`f*&|538R^&?$eR<(Kzi;yUSH=*%y(l71N7LdUxs@>|XK4c|os^d~_N z^UWnde-8Ap{?QHkGoXi!&vDSVT@(KPBVaNTB zgWmrBaNpP@=-WUK>p#<=Uk*CVJg0m{T<0w4dqMw^Nw56Aoqr2X>-N1R>|f9?2R&^3 zB|zT`dRYB;%lm4rq1P9NZ)#lV=a~63_C6*-zXkG-GwFu!ep={Q;yC5!ur)P5%z}P5 z-Vf71R-EOB_3sGi&y^}a=!-v4qWtoH*!UeUL!T6SSpKI$e+K!>o8_-$=PlLwEa+z= zf1&*!1G61;OcPH18FA|o(6@meR{sg1-(bF9`9ZsXb%XvS-cP#d37b9+`r_-t0IQu~lTi?=oU80A{Pk?@T33@l^dqEHL z&5e`%H-_u?$ujh5p?}mYU!@awIt%*J<>B%zWQ^!Alb!l8;!Y!=pI+kq1nAG8eZ%^9 zx6s4tcO3LvJ{VrVlc3)Xdf5BZLJw=7SC4;}x=*yVUXj2;1fZwdKy z{+sAw`l(ywhtbDDr~HS}Cqdr^dRY5T%llJi{Z`&(>z7&3pTqlM`U@9sDSBA_MnK>8 zq42%21nB*shuXiOUk-X${~4F}&o=YF%jqAJpzp=|VfN`X=nsOPHs7z@=9C}w#n*?Y zkBh^s=kPwvBxigWagh3wXWIdtR9T|($#?bj{x!^YRRynm4?--sQU1ic;UFEQyx zY|FIJ!+fW+p!ehbu>NDA;V%b0Z2U#a&=W!rtG{m0&;D?D`s1Ltm!MCAz72GAW2gQN z{V`4HgC3UuS)sq&eBao6x8B0(;Lg-=oqZ{n!|D&IKh4}LUUzf6MuB_eJnGa^o`gMf-5fVXjiu@yl7EP zeZpF#?^vKUJN@sO%DK67D-egU4AOKREvh0P$%@5C)hHfX|NZV{-lMI$W`)B?z5>O? zvWnoleCgcWM@Zfc2H&DhE_t_rul_Z2b5G;9YRcfN{DCcy^0puUduGwx+*_-_cea_o zmsDv+RR-|=5b`c1c~6+~$l1XeB)%oc*X5_+z6#=d%H$)bw!?Qe_?AWH=EjKccP1ZA zE)HLYYRZ zBIpldy?QO>f1g3GyxbN*+*C5xc%Yo(mC2A@$*%HP@G`(0G>cfjED3QL;J6C+MjGh zqa1X%!tcPhn)qrAJ_-Xd`0gQj^&Cr66*2h8CN%h-1Ydn4?#Cf{QG?IZ@2aZMUf?7B zRFyFKR@efmyr+Y2ANYESufybXkBctwJp;aN#Mf=|(W1aV~N+r4UT zZkW>ip_wK*&Ye1*1z)sU%COInXOv+XWE_1I;;W9DoBJ^3?;eA1(K)XCMZq_YbZ0P@ zs-7~^U6iiW36s37;JaxJ$7@$DZgu*6B?V16`5On{#9G{uMdes(^4)0jML?sw6?U(i zo4b$rPB;0!06xnip2xtmWWDet{5+Je1K^3abDVNjm%&4^;6~e@_#(suojW)8GfFRG z@F~BAlm7(xW;V>tZ6&_T4L;9U>;vDk=grMMM10p6e3c)y6Q=x4fG=@0e0|goKQ!B6 zk*gi<2jB87b8|P7yt_?#dtLHo!AEz3cM#u$Cf_O--_fYUf zFM7<}T#n>DXY$d8sZ%do!M7WH?=BlM8pL?#F2H!Z+#TKenRd1y08Q*lb!ikF!e@yM$ zYVcM5%1$=|88ukMPWFf{JKapz(Pj1EdGMv8kGC4KJo9(WFX#e&Rxd6UmQ?=5=? z?136dcXxO>+QGMf6!&XVy5pw2kGslI1mESmq-}39`P^-L3;6Bj@1}HrXz*3O+m-Iw;9K$@+|x#U&l-F~ zHk-2dGT_^GCC4^b9lhG<^G=6Z4jFsFcMtYqZy#6*OVQ{ z9^>jr>;V}yf?~`Brzm?Xh$9-^-weqER zF0#&h@s}4__f#x?bdmK?#R>no$of^)&6STHWnKNEdoh4m-bq%{vXU>pwZiI%O;uRm zt>}2P()vW@g-l=c(u(T0RKL5TLiq7_)rr+_kUuFuud;?LYv^QHuagytA6HseSHAk` zN^4&wP?fJ5>!*14rxn%LR9gR9v0AV_!1h!u`S-OU3U8Z{M1;FL-ykp)xW8*eq6Ca-l?g6wdI6yw-VNgSzW=)FX8Zj`D=K0ZtyVu$QGH9rqMIw~Ej0JbK0-YSIsdPU zYW(xLiZ@u+UsqOtud@2~N~D4{^*doC;i2+Snq~UK(e;&UtVSwpQ2t!`@$@C)fBi8{ zPJZ$CnbmJ|{k^pMEcP+RSsd}gV&RoUotsBt#4{bDg)_Sjne4y#S~o6Yz}_c2#kDU4UpjkGp#T0t0<>4alqM*2&5 zJYlGRxT5--71m7^a2QcPJfr$7E3W$CYb#D#Q@x|29UD4(Di+^DW3>9q6&0VYc%}Z5 zj>_M9%UXX-(b`P;p>!hEYmIa`|Hlki8#uoQ`}Z-CzK>Q|*HpAq;e36`RaT>}PrTrr zb{@Dr9bH9#zwxDOT=b~XU*;-aZmqJcRWJKwg*9A_3ffImR#4hUN2}La4SEcE`?Txt z*H)ZZTYX+d-CHbcaPd_Y*7X(D`zk6v>Qp7me@%70oiDUS%hC(2&7u>kN2~X0{Y&}z z%k!5S_)87^r3U^|1ApEch%_Wrf{8}?mA`A`?>}#D{*v1Nl^TGxr=BNX#=rXSUwHoh zwde0S&)f3_!_e>-~KQ=yo3jYlkn0h*-V&NiEk52!8=eRr*uSqC;PX5;UI3Rz|yqf8X z=?MJi30W_ClWvrDUM;^{C5k_^4loCtL1l# z{PxQ4sQg|fzc^~zF26sQ-#^IjQP{Ai=Xm)&Rel@gceVU(k>6hV9hKj!L<0Aif!QUmg{w}tbgGINJ6ZOdXl;D#h zN6%9a2;L$5-6H=V6<%X`ziazz;rEv7C79S$`+CwnRpG~3Ge1oz$u%PXZNfj}N%unF zLH)K5IQ6SZ&v?C9@EH$IZam`m+N&Q1UV#YlnVuCYCn%13ZdLg4*8O592y8tpa8^gLZf|6O8HD31C3QE&>qF`t(z189jwp*H3d5u8G6%%@Ip3Z*fh)q-0q z-T3)}*Ld)L!6P30U4lnF`1OJ(Jou*s@9^NX_fGZD?ZF=voI+a6=cj^G2#fjrT5t+k zF`ufV=&vQ#j0ZnK@L3Q3CczJQ@V5zWt#ZrXBzTPn?-e}a!7mj&>cOuQJmJAVDR_qm zzgO^X5B_bz^B(-Cf{%Of-v~b8!H-fC!4hlIgTG4fDGy#N__POK37p2))LRp(S!iCN z=cB?;;Tz_2m+&XvmQe7qLOcx$pX781PJ6WUd{*#r!QqNm&nZy6DqX=TPK}<=D20r1 z>Ph#jf>Zd1`FvmGM9)bm_%tEzR11zJ7KMSB&%1zAx)Z|xdLjOg;N2el{aC1weqJu~ zsn)~a08VlaoSRS*P8K;!F<}v2vw`uC3jSHar#gRUK7@u)`4L>_=RZ{gEwL!f#e5dMO!B!hq2L>ZxD_~!M+$c_pMK$=TA5I= zPWLudAxo^q2rDq3`vhO=!5S7q`TSOJ3STjw7byc~iACWm=5w;(mwWIt z1*dQn^H~L4wb!bIf+-e-o-<*95uRvGsK0eTS^Nsd^GP>8sBz&(r%}(p1E>0=@D}r# z6aGn&lM~`gm4UIuy4#b_*9!ii2X6pQ`J7&zP!jGEIsYkmehuTB1>bZ6^UnyrP4H)d zQ$8si#(Wl^$ow@TKQF|%GT4?_6c%GXtASH~#)bb>A^u7DDLlq}j>DcJ@lOfA_7mI) zobvOSr@WsL{7Db~Rl%R};NKVgIS>AG;FNB3Z9-+IP5sxhvKTlj|AK2hwqz+WEU_po z#(b6oC;ohULcx06{ZQ};!L^+CzMA=`1y{opZ#KP#@$U6*IlmBmQuITuO2x9SIEDTJ z`^S@?>jkH98}qqca0;(6pDzO^{njy)P;gX;jSvSXKl-)Au*9P98}nHTT(zUr zn_jE-zqPi{5s~ZX-g5`xR3el5&i?h{}v&}-@tOB(k}Q@JxO$I)lURpF8FU%hhJh*sF3-* z=#9*8NxSQQ@?*iLJnjBo2qgJ4f~(;Vfv-M|@dJV%BZ)jBc*Gw3HObmp&-eyH=`r5bLDMIle-($#_=Mm(|L;~Ec!@ZKaJ>M683Y9XSp981-Ps#kQ z`)z-W@maxj{rs=s*2aVqgJx3C$Kv#}#G()@^Z6`rRX@V7MfC z==!g#XZ$9?cgstk6TJC*357HY9=85{TJgWaig@NJdM#+Z`Cs6-te0D}KTXg-?B|#U z%Fl_G*Z=wog<~@5c9-ik;4in|cK-1e;IBjfnf+iwfvrON2f-)PjB9<{fdUf$)EMKs zpWg%gweWgR%R0J7 zYi?#a)3>_$e<1kG1B`2Z^^z9mulX|LQx-p$2tM(7#`A(dAh`8u#&!KvLP5P2{b%hH zjO%(h4LH?XN0*y_t>E3)Gk%%H&%=UG^fIo?`+LC?dl_FM{M+8b@((=1xNes#1kaz( zIRDqOegr(I{!ciYa3=i7o$WEj`!z77+sDzD^U*K=k@=l5s3&)bj<6e}Hk_e@;7x<#b=gIGRsAj|y&m zo$<>A|C8X8*D=0S@Ll+m$~7td?PS6KD0o!z4>LwRTQLlY-x>u29-aT!tqiXJ6ma67 z-Oc=3{!aw&zKwBR&!??o{%8l|6e&l~oq{KBPpH4OoL>t*eGB8dT~2CcIpckdWB99Q zRPdT_FpfXfbEn`3Zf5+4g1K79k@x=(&x@aVS~ zKUerqSg~WZT=bZ3_tVz}_rv>vQ+6@%F#nr{KLb2S{!<}%%SM(nKEiTzeeM_hz%`6( zJ@cv#DVOLC&Hs?#Gg97XB|k5QBBgqm+{%lTWuC+}f?-LL*6`1Bo&A1mpezeVb2 z%q@Q!cu;wNCj2w!x%sQ#&T^(MV|m*9jh#{OEc47lKdcnO~Q;c`M7$ zOFh(yoH4;?C*1sB5rTNFq7QX_{=49_J2~BBBtNHMq70HB z2d>6NmibQ>{)+{#5&iHQ!5`NA(*JdSj!vgPZcBXoY8o_BY1c_Xm|{b@gsbTc0H;QuPP%_#qPQJT}8 z8BHie*UxntA7T6q;eW;U;CkKwoa(to+FjS%rv#5&$8vOkIHiaA$G^t7=HDuKbb@g` z?>{5>?4^vKV)3&w!*cu;;mFZ|QeKlS)}T5wC|7oDF!0Y9$tWa}o+{CG^C zU72I~(MiVd7CCho*DKmK;+*Ynq`mU9W?y1%{k5-FGT=gUR@LxRtKocYfb zJW*i&$lox2s^FgzeD?i}W4ck#F~iJ1``3)?^7aaDZD$DuB1h461{2al5Ao#=(<2paT0#5#`8S!7i98%9~Mp*vjJ~w|#@Q9~hJt269 zhre?)xLm`8R~~EWI9N&68Us$_V)k>KuFn5Y1)sQqaovtb?PC6p_b|TC;^%h3EvcUq z1z);5I6p1GDcu8CvYeRk@6q@jjMoc(Gs&sM`m2HQYQeuC{P}*ymkWN<7|Wme3gcQ& z<^`X6gmFE-_5ly-53|BQ`w;W%cKmG!|EW;aRIbF0%>OEjpYsGiaG6`qnBcQAPW8O- zT`lJ-=6|)w*?Aeuw?4`Edj)?)@JNdBw+Vh01{&qZ5?tqh0C-S&zbE`t?_fD=WxhN9 zol>r8#`U}B3O@N&#`XBU4!Gu#tf-$!;N-Vl{si+MEyR{9SpN6~<2s+;5`0?vmF|bH zd6$&yA?80;-PF5!N>PAevHLW^xcxr+Zlh0;P2G<7RGgXX9b@SevQB4JtAkToBwpdr@z9uu7^I& zFZze6@NWw~>FIAjA{^r@F~ywAEPiUP1&A86zrvJ`@aIG`iw|_({j#4$0j`)-L^tWST1tfffIlCH(9>sze4cY&2Icw!6);K+kNsX z&d(D+VqCZH>A*?;?DdR4Y4I~Gc=rh7I^9=Z&HUq%AKm}g3SKkL{O5?A9|=A^#(1sZ zE3aWW(S3|-`Cr$#r~kj`TH(K(`QIvXx&@#9GULYyeiLx&&ok1Wb-(?J@OMl7X#AIg zTlcXX-Hva1AE!HWCF44uU4qvPGTtlceg-(D8}+n%;{CznZV0$)-xSNy`TxG)kq!s8P~C#yyR2&+xG+4|F(?ho)G*Ikuxj(lXHc;n*^UO zFjB8&eUQ_gmHu;<;EjS${2lY3E%-&igYq*h{NpmNPZ$2XwVc~nj;_zwPH?)>F*kk{ z@F4krFZ?qPGr#WtbAm^1V*JAvKOguI%b%2f)g<^=1W$a1`8NxG-1Wif)&i&Yow=6z z?Qt#mbU)*75IL)m0pg#27vt*$9~C@uG2^-(ZWBEJKE~f9{HNZ)a%Kk^*W=<{f+xBe zf3@&`Oz>{$|62Z!wVYd-U*mUwgykRTVf3T}xUT|ajTo)A5w<@_`7msh?77q+k*5o7&K_z!GjT+cga zeVp?%xtVb-XS3in7czc+1xvaCxNZwZ_`gpBC%=k!f9m%lXIkXHT3$N)6G#{K&w<+$ z3PFz5GeS7&bE#*YpKAn9+?7y#%bDGp6}(2qweFwC+{$wD((YQ$hXt>Bkma?6SB>D-WN>-wfGfRpOG0t#bT1Hm{9BCccE3|_>q^GoWbt#(U$dNu^z&B=-Y0mp z$oy{+{NDs07yWaF;OBo*%KH%W>wa>B;NvonX#Mc_f+ya={JMT#b{ory<{8)WuNQoJ zKjT`?wBWN}W4y`YXVE^EGbQ=d`MD4{*_quQJM&WFM}L_8Fw4>P|5=eU-pja7_uGO; zwlJ>q|K?9|y4{lQ63&iwJ#b2QQqo;3_&0=q_I#G3^Z(}CCI6DHmfsINsQ*7D{Ie;R zbDGFM_R}n<<}${22;M1pbPwaYKi?|&(r18^Oo#W4uA|PAIgXa(xXr_4BFwnP2DUXTXocIUPOTD@4wVKEwR8w==&U7v~8+ zlV<$&!hf&e2X-=kqTtT}r}p}#w3o(Td{=NjmjWk#OU|us6*(QiRsMIee4XyEHGVVW z#|!`1-ONAn2;;g=CIz3mhVd5RKmN1KKO^;{=a+W~KAUHLJ&x`Yyyi=c>-v0J@QIHw z{wk4w(%-QB$hR5S@;3@T^?t_RDE#*b9(_0Cy9KZK9Lt&gG~;!GzeDhvEaSTW-!1sW z7~?vhFPmaH-O`SyikuGv*L5Y3`gu(A_plsY|JLVOPK}IH-M-repB`rZdJ(o;@B^X` zb$fzvpOJj}T6Z>QbI_-ux8U2hrS zHP8=F^sEpG{FBrCpJjgCpFaycsQ>&#^XHgQ`%&*fu`pd!mQQ*q1c%1no7C&9UsXiyZ!np2-9~XRF z=8+SG|679BT+94#5`4v1B>xvPuJiLQ!SmZ0*X?qf;4^>4_)9H*zAyOf7v22N3EnOJ z{AA%j^FdCxL;9h{uMxcF0hXijZwWsAImVZYoY>#8oC9|;{${~H20W<#|6cf~`V z@Gt%<%du`{{CL4Lf{za}zD@9-3Vz@|#`XAJ_O;-2-wK@CtNUZjugB39f=^3-I8Ed) znr8lqo0(siWnA#-4>A4{i=WR6J|q0P-WL5G%bAjKw?z0`1fSiNP>8P2YXzV5w99w2 z9BCJ=$0{FU`IC3DoC`$$7QrXBF|OO?Yl7z`-DcsxzC-}z&&+lVgm&@AE z{Fcbk@^1%DdPCb4BF6fH@K5Y!IkiIgZ!PEJjGrR-%5O-yKEt?f_kRMe+d^2?&*F!f zKXEnlpDM(9!D}`#{PQ^|B2vJS5DAa)&7aZcw^mLGkTk$U0Z zCHU-n8Lt!kD}qmblX2bOUi6rh_k)bTUHC_V2j%mV!awtO%&*He9m4-V!aq68{3nV0 z`tNYM`7GmS3I0C8Ywlq@F8H$V@;s7wg7FH$qrk}@+tJAQxkBgyuI6i*UzloH?-M-# z0Y=^?a%%pb`8)1oT$gKy;N2oex69`=ztqDUMa~nzgX-<=Gr{;y;7YH`xIRPVJQl+L z3*ny@y>yK5ulyd%nfL_fU)SgDz=QJhd*M%9!u)57oK4^7dX7jvA1(O511CB3ArJF; z$qyv|7c;K)PoLnHluOs=F+XJfncJBESP_;J{JFq2>m-kpG zG}XSJX8d)+-zoU)gN*C?zfJI(5yoFB{4aZgpOSje<-Jt!nZITJB_ijmz=QhL(jN!oX9K5th)RFZ?f%UW{+|l} zcqgaZA?Yr8Qt~;<_{oBQfpDxl@>7iKaq%lPU-NU8lb3oc2>+#mckf~O_X++9!K3#vuKQc{FIdij z%NeI{bKMFqX zX_w~&Pi$d+o&RmWWB>WI#39#Iu%H5PHC`nDt3>|If=_&r<+Ka_py1=18Q1yziRKsm zvqAXZ_T;%A58 z-I7n;j(Y{q-^2VG|FYmSf;Whqlm3g-O?;mDbv+CTK7BRgm?qTo4d5@Af15vVepc#d zOG3fAKa2@(^}6xz3ZD1`<65u2;=fr=Fv2xF}=eY>h0|xN@r?!71H^9wop5?{o-s-QQbaRo5>FM6mt3EP$5;w?yzd}#(0%+PgMn`p(}f+`Vr|qfYbakFraKe!x&CY@y+1d&6U2OBD32nT*+~(W8qOAn zQ|Us%>O6l{d#VFaXSTQ@J&;we)W+0LYBb$n^6lEV@BP|3mQZg?sBI`MqtP#;iDfjG zezC>>Vk}zvl~~MQpV8V(x|sI9mmMq?#!`b+BcA#rN-oQ_=WFO_ZEmo4$e##bPnpUM zq<3V8sc#JzM|yi}d#p?r&7T`CW(%p}K&q#IXb|*NDl?SY(Lc03-Ji;!Qz5aD-4?ZZ zf3}#-pkV_!QoXss9JOkCET!^q^%hWoRAyvgU<|Jqf2F`#bg^zv59fMP#q4mgwkMBH z^NtPm)NSfgTVf5DjATc$jl(0`2Xe)>cf{&4S)`Or?aUSiv;CVoW2xlk{!JS?qp7C8 z^kAkxTSz8T$>fI4VoP=~y}du1%JsHw9WK-j7cLvY_aXHGg^M+D=~BCJCuvIMVZH^W?A2v=7aP%Z>3l90&+dkL$PJ(v>Hb)hhFi8! zM2~6P(AlC35;yP^6*bk5PSCb4$1>Qk#FsT5RMHddEfiJ#O! zdJrnN5RisPy-`ZgVR4rZi^o#urUr+Ksq8?$IF_;t)7b=dwSCii72~|~3A)78{Ls|$~Z^{i8vpWiy z(R+plhl}aKV%zri?XjeP4vWX*&ug!$M{Z$Trt(9SrGlq*ku-7}O{qD0XEs%68{OpZ zFTsN+W)7ZpGN-$YfPmS*D}|}5HQxm*B8gU(nSW#F3V<6sS5jZ=n)mQ^nIK$?&QPId zO84|+hlkrb=36AUK1gorN{v3RjkBO3pqi}CZn7Ng#6xQ}7-XbTbsrci_GJqjRWofw zJ5dF=G=fpF3$0hEQmWW_yuM+*-`MQwPj=R;eAyjSPcgB0K8;b=pY2ceLW4H4(b}Gi zCD#R{23t-~JSrWsZDExsI~&QyRVMCWCY38@2iiyk9`stjr7PxYylzDyyQ05; z{a2AL&x7vFcVS z2U8Kvq?Tx+O&$JDjY6YvO5F!!$sVzXpn6ut$@pjvQQWakxIVJ&>4EK;w63mzj4VI_ zP$wZ%BRLv0j_RWDSQ)xcw`b$XU_RZm^YCWn@Hbog_Gl6-qFAbQ|AQ4zHuZ{BXB$<7 zYH9SPE;7Y1?+6x6b-~*AF81pJmacVCnlWHD#0JtZ8%G9n@EW!GW{_BP?a**>z49Ej z1+}KL@{@_)WmM8!hP4+>Wx(CzCP}4XO}uUs3`d@qv0~Vn-HhC{VJ59T zzNOif`T*LStX|G_Dqk!F&rzf5|3?*F^ zUeeW4Qkq&j*CN#bBLu!{)KFVzu~GiDjuFh}D)z$gUwmMwFY)*69#a?a< zdR9oPnC`_!_*w}wj#AK=_|bdjrFF|?=djEdVJZR6$xh{a?4Tu>ul8YZ&u8K|qD5#b zU}~c2O8YHJPjW{zudx{UrFbqa*ctXtvVNPwJ2R70`7lkvN>2nE8?7=0!!->nmS)kU zbO%@&WJ!=6!Al%}mXByXH^LEyym1Y~ljE~vj!=G+Fz-S}T7XUqSlZTMZv*baEeD_f z!=crJnx9j{0qccIeRE*0wAD zob#445s%cQ*&5{mdqheeuWRF7iyU@2*nrzOxE4#fkUTXOMh2Pf?buOYQP|O@_VGPS zwx9)I`JRH6J8ZXlwMg~&+y;vXGl+A->^WGQ9_BySa%cjw($$9L4*4ZL?zYZW>iGU? z9(%}#?tO)o9bgf%bGa!X(I#bb293{yS8mUDL2iBKZ8Mp-oy9t+-L@BIzOi8sPlW1oDVVtw zbHgoH)E$w%SvsC*6P)K_qj_2d+d+rL4vM??k30bdz5 zt<9!0Dyl&>!;uzHC3d8Ar4ZgE|?zA2wl5 zLd-+V)h}uHY#qY5(2HrmA-LY1Z zY}>QmdtO2({lSZHl2O=XSoq@elM9Nc!s(1+VQ9=9Qi$Og9+LMw){7A<7*-)Av*zmq zEirD!z&N*g#8U-?y+RQn8q1%OKlr!4#UJ}o-x3gpjeuaUyavQM`mlQchDJkB40k{r zN<$09oOCU~+g9G?P<~5;jI*xJgSQwCU)u9+Yfp@1YW1cS=fdB&p*00fnOd`9b1J!3 zz_qQLtW@jy8&;g#w#rJaY2UbVMSE)F>eZb|Y#Fau*{)0{9Fj?_X>VJ(YU|ciU9H+w zT5tQPnmOPSk#{-T`DrH~NF!tQ131f(j78&(1jkd2eC8<`kH+d5iN*5$>7H!gP(O~z zB%>$F*=zEEVmTC1dqix1O6F7%GfyqIf#Bp2E;_hqiXC9(CF*R?lZi5U;L+ zvXU^qDdx)KGbk6t3N5Giy_+#@Zp5yaXxar$K`nGf3kBJ@{bIxt($SY89I)|D8I6=i zS6eQDW>~SZEfvE_EK*~p_gQaO<}zv_2mKRQIkQ_!@6k5ld{`G9z){V`XKxx;tx*t3 zUEhKg)!;~~XZLP6PY@wr7(ie%!XgoqjiwGT*>utr3pGF*8E1{KVr(4689EAVuTLST z{>#FI46mht_(ghVp ztNM#NUWs01cSCXMH)7gGunuh&W3uK0Y7C`DbJ<;OJ=E4ceQ6w>gVzpDyB=O5V@J!M zUA;}9i|6i$C;)}-G!LYA!)UscdtlI1;@xVb(O~CW%g$wxXN)cmmA1(`3k^Me*`A$S zDD+opCKXCdBh88Zv7);Z1=olu7oBXy+_NOjp~MT|IY zG)l2W9Xz80i)8w9V7oc8JzWH~9UVq(V@zQ4&_B(_qY^GMMlSzWY<#IODRc`uMdPYj zh=NZBM-XxB5$%dTB4Zo#is&T8;vW4AOP)F>oi23xGHrP|iXLT&VN#(pgrxWJr>dq* zrjYGTjvDH;HpUv2@_slohNDVkql%u%!kgZZ8bH57+*JmL0@H&%Y7DXRqq!*7pf#=T z_;3@&8Zx=wUYsTcyE;*rB%R7Z@xHoLLP}RclIo3|PHIp?o20t(j=o?GT|%8uF3Tif zw1_=i=kc1(mu^pcc*UA1m$LHNK0*gk+myW#i}PP8RAw=ojANM4Xp7akc~z)7hAq)3 z%)N-{Z5qXnG64Z&cr@38;6Kz3MUaBtoWhI|O*w;K5orfiO~A!iqdkS__QKna*3cqR z)ls7xY{Ibaz+j-RK*ueDmMa4^yCr)@2e>`V!5cI!lavJwVW$qDRjBA{!qXGAGvR3s z*Qhl?J326Ffmlml+K6};Ppr~;It;#Pa7fJ>s&!n4hD7Vd;!t~}EpW0Pnoq^4MM=%H z<3x{6P^lc++tYh7VoK;9=CL8LEd)&%rQ?OrJdKFv7Nu=;?LR|J@^rz0&^!dn#E%?) z^Tibtf;>P!ms0kz=1w{ejsQ@c?Zgby1``ZMyHRcNRQpiRPMk8srr8(`Iz3=L1q_Lh z2()_fs?YYqyiK++yZl2U8W8J@DS+%tW!x>i_Ru*gI#v!{+NwwQ`}$%(Mc48X^;B_H zOsSI?RE-AZ+`5XCp4YZSU_AtGBai`uc|3F*p~>UIb1z1s$pEX}Ywf5lAj419;9fp7 zZpWnt*kMsiJsRA)UV@G3QQA%F$g#%dlS9fBBCSeq$Lh3h;MJH4>GIbkO=j*H*`-_6 zy#ZP|2U>l|h0+rK=3G0ZW7SEo7#+CW=%|;eB?l|F{K(+pukAHYVLTd38B?~cpe6@j zzFRO6u2&ww4q9%tVzo7_!%}Ghqm?14P*)7AvoaH-iRuL=JBka11bdB5`V9Jbek_$H z-N|N22bv52pHzVr7;6}5n=nt%t`$vRhFb6#l}=OXCZcsmvs1v1^{Q>u1PJAiJ1WX8 z0qRMk*h|P*iZeSfGzxP7TbQIMP^3nzQ!+yXav;=RyV6=D*{t6|?7LcWV$^q~P*_-0 zO{7XtrY;<=b~5kx94w^RuKAzzOknfIV|>=Q6W4RVC$PFOG(g%{86jSqCd~e!>8=&A zGMMt}RR^W)(!?(1RlrM5r#T0|Jhs+QOp0o)L)3Z!iiGQ{ zej&#mPFKy%8;Xd^k8o0RXm0yE{92BlOZ?AXAidHfM<3v+U3lpP5%&4HmKTRHm!Tnw8xE9{TZIMvPplE4v6fXBf_8LUW$b7dl1< z27;|&>UOkB9692RK45cw=e8$BQt6w6S}r!?fP^3zv>gS@_~IsG*dj<>12{ zbxe~kyuYJqCF!zz$l{=7BqEK5vxU)YQW>D+vkzJTD04aq_oXZjl%f4+v5jzW8962Y z8`(zm%f?Y#3TI)Z3%YIALv4MZlHSPY*6qF!9EIoAazRbZ>vo!`4l0sW-HcNv`M|0A zaOZ(?UUCAAMb!m@YyjEDhNG3WvsEAeA_wcRM+IyB*y7zg(JLQYFVU?m1HPS;p&|An zHp$irZJdCMPXKsGFayRk$=*-_XrqhvzZ|MU3&0iwH@pU$-xMK-d#ak zjx>++PN-*g;an!o%^zB4J9|EXt3}yFOLx*da1fhl_RMhPcJRZDYmYarsb{2sZ4VeA z>w8ci1ERCa^hv+nMBXUa+q`AMrk82D9;6B_z3QFM{1`j9(QC|o+CawAwAfb|+6Aw< z>|@W{hSGKKj8Z)UJwDD@G&kHegwqqF*?!xubn9MH5al+5NS~|SCi_(}hBr!sK!#1C zO@IJDLCGo2zsnZpUMU$FP+|z{Dzit9Z5ENih)a@uN-29u7b5;a7!VV2;eA&p?yDvM zC6)8RtxhBvY)z(M9a2;sZQU*~1uT9jD4>a5h43B*Y1CM##&;}fXlswIJUr^2E(KYj z1;rJG0GV`ZJ1+PgOd)Kir_b0;Mp<38=u$6gi&0gM=T!f>=ao(Ncb+T$bN-KKo z>OqH~mzIq8bQGP}ms_V7gb-2sa+}h#KdyucoBflBSfd_%5SP{X`=wTYaC?w8rI;J7 z6~g`R3mv>NS*n?4gHw09f7-Jv?VFHS!aTd{>h}xI0ILq_%@9#$i`}Rib8~^V*aMBG zn12I1N^xpy)d?bWRH(lG>>xWNQt{e2U2^Uk&y=Zha~`f75eEe`m)Y};&37^3+RC+O zBhx@u6A~H9)ai^nx_LqZK~G3jy)cD>Lx@UmfROXDd{EIm2YEmY(L9=6_eS0PuC=ba zBfxH?vW5>7S4LSs2XEJ)Ej@8c)`^}q-CcIS zmPRg}S}EO&0{Vgq$nBz+qt1Sks}HCdso8Mz>%Cs(iHjM#l*+C-jGm|lWl`{S`DYIK zsVCE0rHopOruf7(&OaO0xE^enhSp&@r9wZU=bVeyoY~^YSs%1RzOlEL+}2`sly$Vn z#4}6lk}^JanRbJY19){fDILs*_RZpDbSe}#OkkgW)6DB6mk3CV6K z*EH?w`Wyj=mZdpZcKXB2cTa(_#&#;)TwZ&_9DG5M@YhgfdB?r@H>AqhRmqE9Owc>% zV-AJFPywz1k6N@nXTe7dJ=qR)q92Omj@0Vm-rA8mfM~yQSaK?)*@p^aTe=VtH0A52 zTcCIx6+n162f8BE4<2TG@j@kx>s5`JH)uy4Y^sM+?@U<(*#Qh)kNIVam{*jE(-)nY zk+D=}cqqk}nz{NM-*pNeUT2|F)ldlDpmDva>)#!|1WD|)%IS?0_x|hr4k?zN`bti; z-cuwp*P=fxP?N~IKyBW^bhJn7kyG{R98f#HSCcNZ4Q8^tW%q4BZ(Q-JSY3lSpsr)l z{zJDFx6lfqv88r-DXN^T1%DVdSV$d|T`t%Y7De4)!=9p+ zwKDZ2=bPZd1N2B`+kK+cEiI)L#6oOiwe@{)`&D{NvQ&B5u#i@lL z=khC}T|LW~L)@`5Rw&#*kwQu1E9SBwb<0tST>JRGUUX2~05A|KO!g33oGuhbbP<`D z_SL|oOOVfCFAE4|RhRc@LtK5CN;GYYy*)weQjeYH8xIY96S4A}=XTo>2c@e%F6TKs zdysZrdwqV|lyW98tqJ{}q4Gx?8VB*Il^nim=JmG>%Qw*$da{9TV1^Zy>czLwyw;GR z>3C)!ZAa%Ob=mB~OVLXAnNXE&?p@LaOdK;{3;(m6$}!fQR7wg62(rGKR*vN8C~+*# z5wsgx+ng_X&@^RVPiMN$Ir@hXWi2V&uaQsOsc;}M1{-!1hDP#+liirELd{X#t8@X} z9;z1fTKExYeT`+-I{Xh^$hAf*Ve^(#M`0U30m0X0Rn&Y}OMaIMMEgEmGp{q*Q1I~E zgny!jKYFRYuF0QQ?9Xi%HC9sHmL3HQnJazQ_LN!og)~g*$X2x`Q85ZiQ+qDc&_lkn zDB2ST$FS|_M9SmJR0^lWd2W2}bTmeXx_kPeT=?6B>WHY%Ex5o{pjb<2J=6mk#HFw0 z@^-C5KKk-JBut-Uftl*DFi2ypgY9Zx-lhDly@hNx#Mern?^IU1p*e!q!oDdn#6+~+ z08kvH7M1!K$R+VnSC^;TbA~#ey#y+)N0tz_B0qvFxTbs|+X?6%ZhIEu3Z*bc$7CVG?{#N@CzRHuxllj)VVrX zL`&LKI;P?aT}O=QlTsLK$QwE{mDq2e1%>v~WeBO22qG?wbyB0r(hG%jj=ONl;yC(a zgv_W<>C#0Byf}8*rnUTuEGGxvqYbpyqmMw_PG);3=~=(teTLG}(xe(o`Egwlb)_c; zE!4^9Ja15m#mytTSXJPD0(~AY)S#vNcjNqt5%q!UR1qJ-wQpykE7J9)yPos8w1hZZ z`=6e8bVr*G(n0!pEg_x3Z+Hpli=X&hO3;lb`~{UHZj9w|IB$o(ZXR#nh(|H|4)xG& z%XFK9TwYAeZsYo%4z<F|XaG97r13aw!m+<1|Wy-P$dPC+5?68_$v{Ftp~z29Wf z<0&*$$lUu*K0nby$%ghEK95cHY!~UsO_^r61vbZ$ zTwSk>kH+fT)|@LnO2si9$rfAakf5*f;yeU~MX#E!Y)4)hFNk|A8Rj%!EE4e5g3@b| z5-W1|>emjGb4+ivlYNH>1xxL~Em_WKUUE|eUoz6J>XvI%Hj|Z)c8) zg|77t1?LSlc5dRS$8Bm^yGNBX1Ke({`56Ls8MM}navP~z9OUCl2uDZTb?C$6$%|26 zA=|-5sYq-4cZ8`cm7m)W%E{>paPr6xzD}a_NLs+Eo?Jo79TfHv4a~H3$f|?aOmTS^-ZNT`(qEu+< zCN7~Y)fwa!_APG|-a&&sKpzK39(^AZCBL2;V)OEelrCwh z&K&4G;Ui|`!ayHz?UAqtgFeKquOA0T(0O>zB8|ST;jGc@x;b*@5c&+JDf>vEFW$>( zzj<}0QdP{~Fj9AU*dZo|xYI9rQ^!k)7HB}eV4s$coo8vur$^N*uk884PDEU+ZEasDiOz4Js&%kswu6}OUzxy zMLE>()mSy7>gH8vSlrRkD~yMYUZ*eeM|<+peLHcz6-5Sl9L999wl^BN2d*AW0sgo< z20l|A-1)Mh;rkgc+#|!oqIaQt29o^L=9n$urnzox=Bvk2qJw0aR zAf^km3yOW`*?rFR<7NYM4Xi#ZhZNxh7SE@0Q{e_jV{PU0PuU&QXE@C-bk29Cf==R; zxI?Ycu!Ur7mGfF2)XGNr2%Ilwr~qH=AF@l@4&R^8s5a|X?k&~h&eI~W`8u8JSus6_ zTw8m@<`mIvW!x+1w#d7>vI^Y@CW}4@M)#CCBS=q1dfn>UahHuU(l)HImHTlMtY4Lf zq5k1y$LdX;L!GIlXN7~nVg>P<#Nl0d69HOiS?x}Ycb5Jz?h_k)+l}JZ5ZC} z9&Fs;$R_T9#qJ(7soIsYgE4UBrQ3=vXCJ!)ze+1K*Y!M}U3{&WQM$lVJDptD89CvN zORX_JI4;ySgk&_-_YV!_@nNi3l43@&(W5WcFL!Hyz3YX` zh)zyJ)0iVX?uGg2N-^eD-IzX+`Kk@&^^#iwokloRL%{9C<(-8Z>Qd zwc{9j^K?!JTrOqE@c5|NZ|%LZx-N?AsG*gz>3mLw&B#?%N=wmq4{gu4F)cWOf8^wG zyRdaZ#&eYu`K zp!(7wZLT;YDxg836GD6B;XA?^1 za|}>RyF$(uqS0u7=*u~RPDZtriCLr`qdu^J7DK$F5#=(EAQ!vZWym)2*;4q@ zHlmxC8lJlRZX1BCSsF05`qV3~z|T6wqPW^+u*i|qrHvP@AMFe~1J*~<>Ze7%rMqHy zY_La%IVctA{TL5r*U_5#{h)ZCC`GQY(@qqclMw;%x;kutMRV7=tMu28%mn7C+8eS1 z`Qn(PeYXc|+0ae44I=|;pG_{jqKQg^l5~}YcW*~C>e|rdbggRK!S@|7w9A}z*Eui@ z&hlG$hhK~N7Q$7*r!1r7n4OiJj`a&Gy6pv8@;>l(j>{bh?QR4K@t@Q$0l# zVJ_Y^RLJO=ejXddhrf5BbIax|bt}(O3rbq$d>}oD8$$6#tI~T8`1%S;MOBbeC%JJO zE8RJR$qTMgXTFfn1|YdWV@eM@_jug|e+;?o!j2`nHb+a}I<_D$gH}Vxb;+%Z=c9$% z;4i0G4%smZ*Mz$2(5SHb>`BoZp@d6YAtR%2t_sv1+TxJeu_?8BWKiu6$bsNOnTy}) z8eF78-C=ZSXO{9~b^*?u+AX*Jc^+SqlP#@Af@Tl;+RX+#Sj3fx`hbzT7FdR&ts&+) zsgXO^nbPJFN^{5qC{DV1aHyD4(m&OD_soCG(cQ zT|j#Swi#ZHLb|88qOb$kz-2aNFBzr}TC}GJwrA2dDzjgrmgFD@J~G&kkEZou^B31#+hN9aacY$CNyjC29?^r;Kb+ZQJ6ivs%?vUR zF(kC!^Gq%>qH$G5nugV-_&^|yse!sU$EwBf{YAPWZ`dFw$M9Ac79F9%TbXo+>)-?N zyN`>}Hu-R_{E0DlbP>6mg94HQsx=^b0~!HtTRPDS*JoXs1J5Pzz`#LgHN80L+P$+* z-8PJV*_kb_>K~z9RdqA^^W$x$Lq4cnbQ3sK(SVrOb`{cjCz2ZOSTb(JzOe_B9uvN! z$rySVJ+`9_-GKfcl~1mE<~^91VN)F~O7)HLDtG4gkyCBRhwA|Uyq&G3O~(@956*h3 z^c)`ewKDr4v@w$mBCMM}7R;J7rVb9#b&@!8Wq(6h-KybUhKl(~^`!h+G;QgXr*pAo znAKl8?_rJNzQ3-~b32l;6;SFk$x1mZZ;R(!TdYcSUvaf@xTd4>)846a=j!cOHM@}q zkMH4R!RX=cvjbks9NnypVS@3`9!inuid;N&-C|ttz;oB3Yj3g)`QUlj>JVv^T_3ydvZm&7YEhtRuc27FN`4nHQ3q)$MNmr z2(H?0Tj$@IEM548Y$DL~N+IrDgUGpgOK4C3Z~>eCU6hqo>V^XBlH>Es6fR7W2z32I z9l8vz+NX$s9O|Th3>I;3dZfQd`&E3WL2i3L%wK&DkndK|rmLss?5Qrm%g0?-_HN<- z2t3&0G^HT{ru0D_bFuY+S9!tyM!ErR1wBf2qHedW(qD9U3|UTtJ8(sGe}4+IK7Z~e zOap6KZbzp(f`@nc7u=JT{jZ>X)3E9F>>4FJz2D{3GpZ<=k1pu*5>usqfM zNqmx;%4@f?w!j+l=}epcV2*eFYk&6I^{Ui)3xzbOoV2Y|U&;MH7TJJW!>%SyX2A*^ z&f>l;`V92Ix~;=`j50D=Rp3Lw6_ceEK?!3@ZRb<6;xZqyzpGC3x@<}qw}X&llAMy} zN((1a=P}FByiHvcUZ${KUUKAUI|J8t_3ZSzxV&$gaYOSDQ*?AUsW3e`ps2=8otymq z6)mHcwKG@G(-tdr;5z){$O}Ai?c}aheFeoWI9dfwT~Z%}S+3N4v`lwe!#D0*dT`{p zKu!U%Uk-ig%pZru)H8%A2h$~sVw5zx?9Az!@?3_;mnpxatFw%^kGs0lvB&{nIEaGu zdU#|V?XHxu=7M5r5M8TAxA7tJg|;I38F#Fb$oo(l~I;Ty8cV+8`cct@mraqQx8#p>bVC>@VrN9~~4?iIJ41JYVLXJrg|m$(v(JI8ZS?_4K?OrpJ2>>H{6z`0|d z3BZW5Z9`f|>Z@UCCk-0__PxY@ah|xeSbb*}k}|mN+NYNtE*bwy^0pD#bdSA6`EsQd0G^ zmhId{I&&9YwudeGbf=-Tjj8>xL|kGl!N=!uN7~8_ZNS=%_Wc~?2JgLIx|ZIe^__O0 zO4jQ#*|9*kO(yelMUHtUDcPfu4f+`@KB{&GV&jooQIhI`Ax9=7P)UingS^avdt*aiatuLk(<)u#Q1K0cx9d zQ2FpE9m++caNjAK}>gkiqG!aA(SduYP~ z_p2Ui#VMDn8XT$Z!^6(Mpv6f_S>eX#jZ5%35?SRQrY^!F>UzvR1J@2)n>mD z7kuP`L2~?Q0nBQfa*nH#l%Z&llFpJ?{myJ*Fx#KfpAjO{#qPewaPm(gfuP2 z?21**;@zXPsZ#AoOC56z*xT1WuQFj)=B}_eHNPhV*~8#RXua|QQ&HJ#ZMVFAIuOB4 zv=wnsU(gIaxsUe1zO?I$)yC~5WMW`&VP&9Be%h+vU9^SSqIz8)d|pz%C+!S589a>} zJ2w||{jqqOJ`SACte|-2q&kb>a=dB>E%{b?&Fb$nL9y~`tD5h|O!_{2uLp48RN{QH z*tz!G*XyVsmhn@muNA6IFW)zRw1ZhWYT7z_{6;||99r14#{vQWqWD2tDj-1^32Lhh z5k5QFHoysz7*G{lZTlH$4@$)t;;m*)m35vb*;iJMYK03Mj&FKIs%Vs`o zPK(OH1H(BiQE`Et$!Tvx#d?ihXU(8)utZ)vb$ejy zUZHbrsvuTco(w{<1Ye?Z@Kxd{HAij9-kymVjG)tbiq%!&@jxkRQ_Az9YD`tQ$dWps9b(2Gj-)|S(hYQQ*TVR2 zSyz4N1dYBC(n|d{6P^r4c#_% z(NUax?H${UPqeeWuyO>SQOvecpzvV2U)fT+-F%dWiEXioJ-5JF0E%DT?M?aM z6$m^CLj|03vCrci=9GNrunc=$#+hAY^mkok?^v`KPBG5*8uh^q!<}q?fg@m5ME5VX zz`d%Cb7yVjO=-cS9@C#*6qTH!U56@deW`|-X=*xil~!1K&UL#>^is-iWjh(N*XT%P zz7`qY)>lHAaoFDK3XjI71)g~ebm%21p|@wLa97hN_uOUT9+*L%=F*K>Dl$sE(cGlI zLYs33Dhky)&wzWfdVJF66TC|=I?Ce6t#s%2Nj*01!XwrW&ZyCz=fmK54h(3kPj#(Y z)wXr8*was5ORNQWqdZ;MpRi8#2WfI0T|seZ5un+q&9Yj*wAtL zs;ui;n}c@A1Z=|iI+!b^IKD$EVR%D@4MW9TZw}W$^TjsmW@h(yDw7?fI5&?N@QW|* zP@4-*FVtoi9Y^aTkE_;y<)_EMod;^?L-UW50|wh@)M^x3w>!h8F~Wp<;}BXTBSZo9-tT-l!JInx_>0=>>pFxxJ+c7w6R?;G8<=7gW66xdibX*1A zi^#Qu$h|PZ(p>Krx{W?(hBC&0c1)BxX6&++pNA>gu%LtQ(~S9MQ>SA)9t8GQAA~s+ z=DFkWSlWG?rv>1cL?xHs>S(uZ9ePKQ_6)|vGq@TbNP?n9AEhyi&{|yOl%}Xr+rLTf zv*3IE8#1}x-fV$_OqG*2NiHt@e=*)zfPZahAjLO=Il3hC(t7 zL)*l1jZ+l^1qZnrt9D0dP~IwCv{FNVVQoUSp{8>?MLn^EDrLt%Gk=Y2tMbx%*$($E zsg%uG!5C&VINb{FP&;aA9S9RreWz3hTJi_?;C4;nlAYwH6?LigTJO#y$iVAJK;FIP zmD)=)KaAko9SGFs$~Vfo%h$lc9{*^Dzut~Jl9L@dcV|3;%Mm)Jw<2>fozKjiOPAte z2&;qn)W$fPI?9{O1NH_p2Kxfll&y_Jf>pv}fk?seekZOt=ubn@QuKP8L^;wWY5Z;b zfW=`+=g2&-bih#dmKh39kLMuoo!YjvJI0Y~UMvW5F=|Usgy@Bh^5y$)PuS48%IDu#<{@gioUimC?gO=h zxPuhfe6N178&Q`jPNzNT8p-xFE;s3W{;!c#HDgRrK?e#Q9i>}YTb^!bm$Qon0|kSs zo|fFs7u>tbUDhLJ%fc>pDH#O<^&T1p0!pAg;2o59TS+B^t*@|+ZEE99u$U<9N$F&# zq6c#Gzxz0_UNP!QeDE&jvm2p#SEV{vtCbRp=84gp?}2~se1`@$Y<}-LOm`QTha*TzoxQM7rU-BasNDLDJw1LD>oWoUK{>TKE-iPWB44gnmmnZlXePdNFMI<0HZKYc)ki5km>*yhfu2P=Tsb}tYD3nbs7haOQX*D zf&>R$nyUq;L+B1nVSn`9d2er_I*!K+2t^&uw&FNAIgzuiYS;~v5pGW@rYh5hFoXph zUk%1Vo!1$&(C z<&&1|d^FL%8Pst}Uo74oC@JF#SvsZ1zJI2WDR5L&^51s4xTj?+ysAb99ry zMW%|{h}Cg;6=qk8{^GO$_Von#V3^tn5!b1Fjmp$B)N}0pTz?>Jfa)TN)Qvcngj@v& z9R`by7&I*FB1HV)WOaxxGIrU#UyxK2FDwlTDIo=il`xt!c2HN&A`A6B7=dP>IrTUi zFs@#L4!arlIRd2dx!xkKEwp#&J@bb8#0##=!eC~%f_Z+?9L2`WYtP6MPT-p816-ZZ z;OD03#H0jlV#O-M(t96*_c9l=;gvN5Z349~b{wgrlHbJg&H`Z_lQwcZ{;GK{j_q0r zt2XTjgYI%E11^nz-7Ke?egc=NzHc^#sCH!oa!jr(O(ofPD1f<*n=iE)6MU(&{eBzF zv9&|!S*0x`7O%as9#@(_d%DX7>PP2ow=L+@!@}H^G_g=*3+I*3AbXo6eB-lY6mhd^ zo?=(G{&F_$7}XE3S1D*~R1I7^nfa-UIKIvhbeTW*UT68{-NRA-`!b(EbFT~OQ`+u{ zrh%hmlohCIipp$yP1Tv1dLvGTBW$Rgb}GKl0=VbTM+9=tzhwzYZI^T4@NQ+?~&iqU8Q@-td`!9Sg z<6H2*k4CHPzvld!^L?*B;5zTWi1;x)vS`U&#;3E!QLN_3|U{ANn$mpM7Vhr?4C za#ZfOJ`**yKBFou`M~T?p67ptGU2aSc_@phR#{kJMLsS>_;Yx2{v~44Lf)@RBf2&v z_cIlK7RFw{NcqC~qTv-B++BVS>eo3xAGNZWH=y;#G*@!;-@-`yFKw=MzrJynh5U?LtkN}{paF>rskqVA=Sk)H}AiX z@NqnRO}+cZ_pTeCpO^iqkD}~E!uOr;|0jPIymVs!f%60BC##ErB1oa@c$oz(MI)ekI>z%4M>m_F9QEal?>;uZ zdm9M|=i~(B=tq=EjPv)L?>XPJZCD0^s;^bdtac)Pw2%~mmw3_o5?^wl| zfACV3y&@&9IUkp7g97R!9-g0nf92m3&knPa=kKlmhaXi)4U3Lrr(phxkMC@%dl`_q F;TN_`NNfNA From 78e7ae65e2e955f7fc4797bae399f6ec173e2d6a Mon Sep 17 00:00:00 2001 From: Lorenc Bushi Date: Sun, 19 Jan 2025 21:58:45 -0500 Subject: [PATCH 18/53] Improve L0_interop_test.cpp to avoid false positives --- sycl/test-e2e/DeviceImageBackendContent/L0_interop_test.cpp | 2 ++ 1 file changed, 2 insertions(+) diff --git a/sycl/test-e2e/DeviceImageBackendContent/L0_interop_test.cpp b/sycl/test-e2e/DeviceImageBackendContent/L0_interop_test.cpp index fc6cfeed3e3a4..ab0b920d5626d 100644 --- a/sycl/test-e2e/DeviceImageBackendContent/L0_interop_test.cpp +++ b/sycl/test-e2e/DeviceImageBackendContent/L0_interop_test.cpp @@ -34,6 +34,7 @@ int main() { sycl::kernel_id iota_id = syclexp::get_kernel_id(); sycl::kernel k_iota = Bundle.get_kernel(iota_id); int *ptr = sycl::malloc_shared(1, q); + *ptr = 0; q.submit([&](sycl::handler &cgh) { cgh.set_args(ptr); cgh.parallel_for(sycl::nd_range{{1}, {1}}, k_iota); @@ -84,6 +85,7 @@ int main() { ZeKernel}, ctxt); int *ptr_twin = sycl::malloc_shared(1, q); + *ptr_twin = 1; q.submit([&](sycl::handler &cgh) { cgh.set_args(ptr_twin); cgh.parallel_for(sycl::nd_range{{1}, {1}}, k_iota_twin); From b2594e8ad6f50d77cb1b3e14bff56f4d6c763f1b Mon Sep 17 00:00:00 2001 From: Lorenc Bushi Date: Tue, 21 Jan 2025 10:31:56 -0500 Subject: [PATCH 19/53] Use __has_include to guard the inclusion of Co-authored-by: John Pennycook --- sycl/include/sycl/kernel_bundle.hpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/sycl/include/sycl/kernel_bundle.hpp b/sycl/include/sycl/kernel_bundle.hpp index c3060c7c96f73..249720f53f958 100644 --- a/sycl/include/sycl/kernel_bundle.hpp +++ b/sycl/include/sycl/kernel_bundle.hpp @@ -32,7 +32,7 @@ #include // for function #include // for distance #include // for shared_ptr, operator==, hash -#ifdef __cpp_lib_span +#if __has_include() #include // for span #endif #include // for string From 5b8876182de6548e87d3af77a756259a6063ec87 Mon Sep 17 00:00:00 2001 From: Lorenc Bushi Date: Tue, 21 Jan 2025 11:57:04 -0500 Subject: [PATCH 20/53] Update basic_test.cpp --- sycl/test-e2e/DeviceImageBackendContent/basic_test.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/sycl/test-e2e/DeviceImageBackendContent/basic_test.cpp b/sycl/test-e2e/DeviceImageBackendContent/basic_test.cpp index 2e33f9c610d49..c2e3f2c330b00 100644 --- a/sycl/test-e2e/DeviceImageBackendContent/basic_test.cpp +++ b/sycl/test-e2e/DeviceImageBackendContent/basic_test.cpp @@ -2,7 +2,7 @@ // RUN: %{build} -o %t.out // RUN: %{run} %t.out -#ifdef __cpp_lib_span +#if __has_include() #include #endif #include From ccdd15cd748459c87fb6e789ed07e653491ef75d Mon Sep 17 00:00:00 2001 From: Lorenc Bushi Date: Wed, 22 Jan 2025 12:08:23 -0800 Subject: [PATCH 21/53] Apply feedback --- sycl/include/sycl/kernel_bundle.hpp | 11 +++-- sycl/source/kernel_bundle.cpp | 13 +++-- .../DeviceImageBackendContent/basic_test.cpp | 17 ------- sycl/test/abi/sycl_symbols_linux.dump | 5 +- .../negative_test.cpp | 47 +++++++++++++++++++ 5 files changed, 62 insertions(+), 31 deletions(-) create mode 100644 sycl/test/extensions/DeviceImageBackendContent/negative_test.cpp diff --git a/sycl/include/sycl/kernel_bundle.hpp b/sycl/include/sycl/kernel_bundle.hpp index 249720f53f958..fc1805a65e290 100644 --- a/sycl/include/sycl/kernel_bundle.hpp +++ b/sycl/include/sycl/kernel_bundle.hpp @@ -33,7 +33,7 @@ #include // for distance #include // for shared_ptr, operator==, hash #if __has_include() -#include // for span +#include #endif #include // for string #include // for enable_if_t, remove_refer... @@ -122,9 +122,8 @@ class __SYCL_EXPORT device_image_plain { #if (!defined(_HAS_STD_BYTE) || _HAS_STD_BYTE != 0) std::vector ext_oneapi_get_backend_content() const; -#ifdef __cpp_lib_span - std::span ext_oneapi_get_backend_content_view() const; -#endif // __cpp_lib_span + std::pair + ext_oneapi_get_backend_content_view() const; #endif // HAS_STD_BYTE @@ -174,7 +173,9 @@ class device_image : public detail::device_image_plain, template > std::span ext_oneapi_get_content_backend_view() const { - return device_image_plain::ext_oneapi_get_backend_content_view(); + return std::span{ + device_image_plain::ext_oneapi_get_backend_content_view().first, + device_image_plain::ext_oneapi_get_backend_content_view().second}; } #endif // __cpp_lib_span #endif // _HAS_STD_BYTE diff --git a/sycl/source/kernel_bundle.cpp b/sycl/source/kernel_bundle.cpp index 3e987de513bb4..bdaea463cf3cb 100644 --- a/sycl/source/kernel_bundle.cpp +++ b/sycl/source/kernel_bundle.cpp @@ -57,15 +57,14 @@ device_image_plain::ext_oneapi_get_backend_content() const { impl->get_bin_image_ref()->getRawData().BinaryEnd)); } -#ifdef __cpp_lib_span -std::span +std::pair device_image_plain::ext_oneapi_get_backend_content_view() const { - return std::span(reinterpret_cast( - impl->get_bin_image_ref()->getRawData().BinaryStart), - reinterpret_cast( - impl->get_bin_image_ref()->getRawData().BinaryEnd)); + return std::make_pair( + reinterpret_cast( + impl->get_bin_image_ref()->getRawData().BinaryStart), + reinterpret_cast( + impl->get_bin_image_ref()->getRawData().BinaryEnd)); } -#endif //////////////////////////// ///// kernel_bundle_plain diff --git a/sycl/test-e2e/DeviceImageBackendContent/basic_test.cpp b/sycl/test-e2e/DeviceImageBackendContent/basic_test.cpp index c2e3f2c330b00..da87d3d1edff9 100644 --- a/sycl/test-e2e/DeviceImageBackendContent/basic_test.cpp +++ b/sycl/test-e2e/DeviceImageBackendContent/basic_test.cpp @@ -1,4 +1,3 @@ -// RUN: %{build} -fsyntax-only -DTEST_API_VIOLATION=1 -Xclang -verify -Xclang -verify-ignore-unexpected=note,warning // RUN: %{build} -o %t.out // RUN: %{run} %t.out @@ -58,21 +57,5 @@ int main() { } #endif } - -#ifdef TEST_API_VIOLATION - // Check that the ext_oneapi_get_backend_content and the - // ext_oneapi_get_backend_content_view of the content functions are not - // available - // when the image is not in the executable state. - - auto input_bundle = - sycl::get_kernel_bundle(ctxt, {id}); - // expected-error@+1 {{no matching member function for call to 'ext_oneapi_get_backend_content'}} - bytes = (*input_bundle.begin()).ext_oneapi_get_backend_content(); -#ifdef _cpp_lib_span - // expected-error@+1 {{no matching member function for call to 'ext_oneapi_get_backend_content_view'}} - bytes_view = (*input_bundle.begin()).ext_oneapi_get_backend_content_view(); -#endif // __cpp_lib_span -#endif // TEST_API_VIOLATION return 0; } diff --git a/sycl/test/abi/sycl_symbols_linux.dump b/sycl/test/abi/sycl_symbols_linux.dump index 2e20339cc202a..c9ef3c269ce57 100644 --- a/sycl/test/abi/sycl_symbols_linux.dump +++ b/sycl/test/abi/sycl_symbols_linux.dump @@ -3524,14 +3524,14 @@ _ZN4sycl3_V17handler20associateWithHandlerEPNS0_6detail30UnsampledImageAccessorB _ZN4sycl3_V17handler20memcpyToDeviceGlobalEPKvS3_bmm _ZN4sycl3_V17handler20setKernelCacheConfigENS1_23StableKernelCacheConfigE _ZN4sycl3_V17handler20setStateSpecConstSetEv -_ZN4sycl3_V17handler21setKernelWorkGroupMemEm _ZN4sycl3_V17handler21setUserFacingNodeTypeENS0_3ext6oneapi12experimental9node_typeE _ZN4sycl3_V17handler22ext_oneapi_fill2d_implEPvmPKvmmm _ZN4sycl3_V17handler22memcpyFromDeviceGlobalEPvPKvbmm _ZN4sycl3_V17handler22setHandlerKernelBundleENS0_6kernelE _ZN4sycl3_V17handler22setHandlerKernelBundleERKSt10shared_ptrINS0_6detail18kernel_bundle_implEE _ZN4sycl3_V17handler22setKernelClusterLaunchENS0_5rangeILi3EEEi -_ZN4sycl3_V17handler22setKernelIsCooperativeEb +_ZN4sycl3_V17handler22setKernelIsCooperativeEbi +_ZN4sycl3_V17handler21setKernelWorkGroupMemEm _ZN4sycl3_V17handler24GetRangeRoundingSettingsERmS2_S2_ _ZN4sycl3_V17handler24ext_intel_read_host_pipeENS0_6detail11string_viewEPvmb _ZN4sycl3_V17handler24ext_oneapi_memcpy2d_implEPvmPKvmmm @@ -3739,6 +3739,7 @@ _ZNK4sycl3_V16detail18device_image_plain10has_kernelERKNS0_9kernel_idE _ZNK4sycl3_V16detail18device_image_plain10has_kernelERKNS0_9kernel_idERKNS0_6deviceE _ZNK4sycl3_V16detail18device_image_plain22ext_oneapi_get_backendEv _ZNK4sycl3_V16detail18device_image_plain30ext_oneapi_get_backend_contentEv +_ZNK4sycl3_V16detail18device_image_plain35ext_oneapi_get_backend_content_viewEv _ZNK4sycl3_V16detail18device_image_plain9getNativeEv _ZNK4sycl3_V16detail19kernel_bundle_plain10get_kernelERKNS0_9kernel_idE _ZNK4sycl3_V16detail19kernel_bundle_plain10has_kernelERKNS0_9kernel_idE diff --git a/sycl/test/extensions/DeviceImageBackendContent/negative_test.cpp b/sycl/test/extensions/DeviceImageBackendContent/negative_test.cpp new file mode 100644 index 0000000000000..e48031f7767ee --- /dev/null +++ b/sycl/test/extensions/DeviceImageBackendContent/negative_test.cpp @@ -0,0 +1,47 @@ +// RUN: %clang -fsycl -fsyntax-only -Xclang -verify -Xclang -verify-ignore-unexpected=note,warning %s + +#include + +class kernel; + +sycl::device d; +sycl::queue q{d}; +sycl::context ctxt = q.get_context(); +sycl::kernel_id id = sycl::get_kernel_id(); + +int main() { + // Check that the ext_oneapi_get_backend_content and the + // ext_oneapi_get_backend_content_view of the content functions are not + // available + // when the image is not in the executable state. + + auto input_bundle = + sycl::get_kernel_bundle(ctxt, {id}); + // expected-error@+1 {{no matching member function for call to 'ext_oneapi_get_backend_content'}} + (*input_bundle.begin()).ext_oneapi_get_backend_content(); +#ifdef __cpp_lib_span + // expected-error@+1 {{no matching member function for call to 'ext_oneapi_get_backend_content_view'}} + (*input_bundle.begin()).ext_oneapi_get_backend_content_view(); +#endif + + auto object_bundle = + sycl::get_kernel_bundle(ctxt, {id}); + // expected-error@+1 {{no matching member function for call to 'ext_oneapi_get_backend_content'}} + (*input_bundle.begin()).ext_oneapi_get_backend_content(); +#ifdef __cpp_lib_span + // expected-error@+1 {{no matching member function for call to 'ext_oneapi_get_backend_content_view'}} + (*input_bundle.begin()).ext_oneapi_get_backend_content_view(); +#endif + + auto source_bundle = + sycl::get_kernel_bundle(ctxt, + {id}); + // expected-error@+1 {{no matching member function for call to 'ext_oneapi_get_backend_content'}} + (*input_bundle.begin()).ext_oneapi_get_backend_content(); +#ifdef __cpp_lib_span + // expected-error@+1 {{no matching member function for call to 'ext_oneapi_get_backend_content_view'}} + (*input_bundle.begin()).ext_oneapi_get_backend_content_view(); +#endif + + return 0; +} From 21f4a9b44ee99e4f7d55e273832f1f7f6732b970 Mon Sep 17 00:00:00 2001 From: Lorenc Bushi Date: Wed, 22 Jan 2025 16:01:33 -0500 Subject: [PATCH 22/53] Update sycl_symbols_linux.dump --- sycl/test/abi/sycl_symbols_linux.dump | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/sycl/test/abi/sycl_symbols_linux.dump b/sycl/test/abi/sycl_symbols_linux.dump index 6a98c2c47f939..12eb791f9da88 100644 --- a/sycl/test/abi/sycl_symbols_linux.dump +++ b/sycl/test/abi/sycl_symbols_linux.dump @@ -3533,7 +3533,7 @@ _ZN4sycl3_V17handler22memcpyFromDeviceGlobalEPvPKvbmm _ZN4sycl3_V17handler22setHandlerKernelBundleENS0_6kernelE _ZN4sycl3_V17handler22setHandlerKernelBundleERKSt10shared_ptrINS0_6detail18kernel_bundle_implEE _ZN4sycl3_V17handler22setKernelClusterLaunchENS0_5rangeILi3EEEi -_ZN4sycl3_V17handler22setKernelIsCooperativeEbi +_ZN4sycl3_V17handler22setKernelIsCooperativeEb _ZN4sycl3_V17handler21setKernelWorkGroupMemEm _ZN4sycl3_V17handler24GetRangeRoundingSettingsERmS2_S2_ _ZN4sycl3_V17handler24ext_intel_read_host_pipeENS0_6detail11string_viewEPvmb From 91853e03a8e9d014e95b4284de7143321144a776 Mon Sep 17 00:00:00 2001 From: Lorenc Bushi Date: Wed, 22 Jan 2025 16:02:25 -0500 Subject: [PATCH 23/53] Update sycl_symbols_linux.dump --- sycl/test/abi/sycl_symbols_linux.dump | 1 - 1 file changed, 1 deletion(-) diff --git a/sycl/test/abi/sycl_symbols_linux.dump b/sycl/test/abi/sycl_symbols_linux.dump index 12eb791f9da88..b295669c0ac91 100644 --- a/sycl/test/abi/sycl_symbols_linux.dump +++ b/sycl/test/abi/sycl_symbols_linux.dump @@ -3534,7 +3534,6 @@ _ZN4sycl3_V17handler22setHandlerKernelBundleENS0_6kernelE _ZN4sycl3_V17handler22setHandlerKernelBundleERKSt10shared_ptrINS0_6detail18kernel_bundle_implEE _ZN4sycl3_V17handler22setKernelClusterLaunchENS0_5rangeILi3EEEi _ZN4sycl3_V17handler22setKernelIsCooperativeEb -_ZN4sycl3_V17handler21setKernelWorkGroupMemEm _ZN4sycl3_V17handler24GetRangeRoundingSettingsERmS2_S2_ _ZN4sycl3_V17handler24ext_intel_read_host_pipeENS0_6detail11string_viewEPvmb _ZN4sycl3_V17handler24ext_oneapi_memcpy2d_implEPvmPKvmmm From 70c168a687027d48bbd3439b47adb04d94498b20 Mon Sep 17 00:00:00 2001 From: Lorenc Bushi Date: Wed, 22 Jan 2025 17:27:43 -0800 Subject: [PATCH 24/53] Add windows symbols --- sycl/test/abi/sycl_symbols_windows.dump | 29 ++++++++++++++----------- 1 file changed, 16 insertions(+), 13 deletions(-) diff --git a/sycl/test/abi/sycl_symbols_windows.dump b/sycl/test/abi/sycl_symbols_windows.dump index 5d09568968d70..a233b2214670d 100644 --- a/sycl/test/abi/sycl_symbols_windows.dump +++ b/sycl/test/abi/sycl_symbols_windows.dump @@ -282,9 +282,9 @@ ??0SampledImageAccessorBaseHost@detail@_V1@sycl@@QEAA@$$QEAV0123@@Z ??0SampledImageAccessorBaseHost@detail@_V1@sycl@@QEAA@AEBV0123@@Z ??0SampledImageAccessorBaseHost@detail@_V1@sycl@@QEAA@V?$range@$02@23@PEAXHHV?$id@$02@23@W4image_channel_type@23@W4image_channel_order@23@Uimage_sampler@23@AEBVproperty_list@23@@Z -??0SubmissionInfo@detail@_V1@sycl@@QEAA@$$QEAV0123@@Z -??0SubmissionInfo@detail@_V1@sycl@@QEAA@AEBV0123@@Z ??0SubmissionInfo@detail@_V1@sycl@@QEAA@XZ +??0SubmissionInfo@detail@_V1@sycl@@QEAA@AEBV0123@@Z +??0SubmissionInfo@detail@_V1@sycl@@QEAA@$$QEAV0123@@Z ??0UnsampledImageAccessorBaseHost@detail@_V1@sycl@@IEAA@AEBV?$shared_ptr@VUnsampledImageAccessorImplHost@detail@_V1@sycl@@@std@@@Z ??0UnsampledImageAccessorBaseHost@detail@_V1@sycl@@QEAA@$$QEAV0123@@Z ??0UnsampledImageAccessorBaseHost@detail@_V1@sycl@@QEAA@AEBV0123@@Z @@ -330,12 +330,18 @@ ??0device_image_plain@detail@_V1@sycl@@QEAA@AEBV?$shared_ptr@Vdevice_image_impl@detail@_V1@sycl@@@std@@@Z ??0device_selector@_V1@sycl@@QEAA@AEBV012@@Z ??0device_selector@_V1@sycl@@QEAA@XZ -??0dynamic_command_group@experimental@oneapi@ext@_V1@sycl@@QEAA@$$QEAV012345@@Z -??0dynamic_command_group@experimental@oneapi@ext@_V1@sycl@@QEAA@AEBV012345@@Z -??0dynamic_command_group@experimental@oneapi@ext@_V1@sycl@@QEAA@AEBV?$command_graph@$0A@@12345@AEBV?$vector@V?$function@$$A6AXAEAVhandler@_V1@sycl@@@Z@std@@V?$allocator@V?$function@$$A6AXAEAVhandler@_V1@sycl@@@Z@std@@@2@@std@@@Z ??0dynamic_parameter_base@detail@experimental@oneapi@ext@_V1@sycl@@QEAA@$$QEAV0123456@@Z ??0dynamic_parameter_base@detail@experimental@oneapi@ext@_V1@sycl@@QEAA@AEBV0123456@@Z ??0dynamic_parameter_base@detail@experimental@oneapi@ext@_V1@sycl@@QEAA@V?$command_graph@$0A@@23456@_KPEBX@Z +?get_active_index@dynamic_command_group@experimental@oneapi@ext@_V1@sycl@@QEBA_KXZ +??4dynamic_command_group@experimental@oneapi@ext@_V1@sycl@@QEAAAEAV012345@$$QEAV012345@@Z +?addImpl@modifiable_command_graph@detail@experimental@oneapi@ext@_V1@sycl@@IEAA?AVnode@34567@AEAVdynamic_command_group@34567@AEBV?$vector@Vnode@experimental@oneapi@ext@_V1@sycl@@V?$allocator@Vnode@experimental@oneapi@ext@_V1@sycl@@@std@@@std@@@Z +??0dynamic_command_group@experimental@oneapi@ext@_V1@sycl@@QEAA@AEBV?$command_graph@$0A@@12345@AEBV?$vector@V?$function@$$A6AXAEAVhandler@_V1@sycl@@@Z@std@@V?$allocator@V?$function@$$A6AXAEAVhandler@_V1@sycl@@@Z@std@@@2@@std@@@Z +??0dynamic_command_group@experimental@oneapi@ext@_V1@sycl@@QEAA@$$QEAV012345@@Z +??0dynamic_command_group@experimental@oneapi@ext@_V1@sycl@@QEAA@AEBV012345@@Z +?set_active_index@dynamic_command_group@experimental@oneapi@ext@_V1@sycl@@QEAAX_K@Z +??1dynamic_command_group@experimental@oneapi@ext@_V1@sycl@@QEAA@XZ +??4dynamic_command_group@experimental@oneapi@ext@_V1@sycl@@QEAAAEAV012345@AEBV012345@@Z ??0event@_V1@sycl@@AEAA@V?$shared_ptr@Vevent_impl@detail@_V1@sycl@@@std@@@Z ??0event@_V1@sycl@@QEAA@$$QEAV012@@Z ??0event@_V1@sycl@@QEAA@AEBV012@@Z @@ -473,7 +479,6 @@ ??1device@_V1@sycl@@QEAA@XZ ??1device_image_plain@detail@_V1@sycl@@QEAA@XZ ??1device_selector@_V1@sycl@@UEAA@XZ -??1dynamic_command_group@experimental@oneapi@ext@_V1@sycl@@QEAA@XZ ??1dynamic_parameter_base@detail@experimental@oneapi@ext@_V1@sycl@@QEAA@XZ ??1event@_V1@sycl@@QEAA@XZ ??1exception@_V1@sycl@@UEAA@XZ @@ -551,8 +556,6 @@ ??4device_image_plain@detail@_V1@sycl@@QEAAAEAV0123@$$QEAV0123@@Z ??4device_image_plain@detail@_V1@sycl@@QEAAAEAV0123@AEBV0123@@Z ??4device_selector@_V1@sycl@@QEAAAEAV012@AEBV012@@Z -??4dynamic_command_group@experimental@oneapi@ext@_V1@sycl@@QEAAAEAV012345@$$QEAV012345@@Z -??4dynamic_command_group@experimental@oneapi@ext@_V1@sycl@@QEAAAEAV012345@AEBV012345@@Z ??4dynamic_parameter_base@detail@experimental@oneapi@ext@_V1@sycl@@QEAAAEAV0123456@$$QEAV0123456@@Z ??4dynamic_parameter_base@detail@experimental@oneapi@ext@_V1@sycl@@QEAAAEAV0123456@AEBV0123456@@Z ??4event@_V1@sycl@@QEAAAEAV012@$$QEAV012@@Z @@ -651,8 +654,8 @@ ?GDBMethodsAnchor@UnsampledImageAccessorBaseHost@detail@_V1@sycl@@IEAAXXZ ?GetRangeRoundingSettings@handler@_V1@sycl@@AEAAXAEA_K00@Z ?HasAssociatedAccessor@handler@_V1@sycl@@AEBA_NPEAVAccessorImplHost@detail@23@W4target@access@23@@Z -?PostProcessorFunc@SubmissionInfo@detail@_V1@sycl@@QEAAAEAV?$optional@V?$function@$$A6AX_N0AEAVevent@_V1@sycl@@@Z@std@@@234@XZ ?PostProcessorFunc@SubmissionInfo@detail@_V1@sycl@@QEBAAEBV?$optional@V?$function@$$A6AX_N0AEAVevent@_V1@sycl@@@Z@std@@@234@XZ +?PostProcessorFunc@SubmissionInfo@detail@_V1@sycl@@QEAAAEAV?$optional@V?$function@$$A6AX_N0AEAVevent@_V1@sycl@@@Z@std@@@234@XZ ?PushBack@exception_list@_V1@sycl@@AEAAX$$QEAVexception_ptr@std@@@Z ?PushBack@exception_list@_V1@sycl@@AEAAXAEBVexception_ptr@std@@@Z ?RangeRoundingTrace@handler@_V1@sycl@@AEAA_NXZ @@ -3704,7 +3707,6 @@ ?addHostAccessorAndWait@detail@_V1@sycl@@YAXPEAVAccessorImplHost@123@@Z ?addHostSampledImageAccessorAndWait@detail@_V1@sycl@@YAXPEAVSampledImageAccessorImplHost@123@@Z ?addHostUnsampledImageAccessorAndWait@detail@_V1@sycl@@YAXPEAVUnsampledImageAccessorImplHost@123@@Z -?addImpl@modifiable_command_graph@detail@experimental@oneapi@ext@_V1@sycl@@IEAA?AVnode@34567@AEAVdynamic_command_group@34567@AEBV?$vector@Vnode@experimental@oneapi@ext@_V1@sycl@@V?$allocator@Vnode@experimental@oneapi@ext@_V1@sycl@@@std@@@std@@@Z ?addImpl@modifiable_command_graph@detail@experimental@oneapi@ext@_V1@sycl@@IEAA?AVnode@34567@AEBV?$vector@Vnode@experimental@oneapi@ext@_V1@sycl@@V?$allocator@Vnode@experimental@oneapi@ext@_V1@sycl@@@std@@@std@@@Z ?addImpl@modifiable_command_graph@detail@experimental@oneapi@ext@_V1@sycl@@IEAA?AVnode@34567@V?$function@$$A6AXAEAVhandler@_V1@sycl@@@Z@std@@AEBV?$vector@Vnode@experimental@oneapi@ext@_V1@sycl@@V?$allocator@Vnode@experimental@oneapi@ext@_V1@sycl@@@std@@@std@@@Z ?addLifetimeSharedPtrStorage@handler@_V1@sycl@@AEAAXV?$shared_ptr@$$CBX@std@@@Z @@ -3842,13 +3844,14 @@ ?ext_oneapi_fill2d_impl@handler@_V1@sycl@@AEAAXPEAX_KPEBX111@Z ?ext_oneapi_get_backend@device_image_plain@detail@_V1@sycl@@QEBA?AW4backend@34@XZ ?ext_oneapi_get_backend_content@device_image_plain@detail@_V1@sycl@@QEBA?AV?$vector@W4byte@std@@V?$allocator@W4byte@std@@@2@@std@@XZ +?ext_oneapi_get_backend_content_view@device_image_plain@detail@_V1@sycl@@QEBA?AU?$pair@PEBW4byte@std@@PEBW412@@std@@XZ ?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 ?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 -?ext_oneapi_get_last_event@queue@_V1@sycl@@QEBA?AV?$optional@Vevent@_V1@sycl@@@std@@XZ ?ext_oneapi_get_last_event_impl@queue@_V1@sycl@@AEBA?AV?$optional@Vevent@_V1@sycl@@@detail@23@XZ +?ext_oneapi_get_last_event@queue@_V1@sycl@@QEBA?AV?$optional@Vevent@_V1@sycl@@@std@@XZ ?ext_oneapi_get_state@queue@_V1@sycl@@QEBA?AW4queue_state@experimental@oneapi@ext@23@XZ ?ext_oneapi_graph@handler@_V1@sycl@@QEAAXV?$command_graph@$00@experimental@oneapi@ext@23@@Z ?ext_oneapi_graph@queue@_V1@sycl@@QEAA?AVevent@23@V?$command_graph@$00@experimental@oneapi@ext@23@AEBUcode_location@detail@23@@Z @@ -4014,7 +4017,6 @@ ?getType@handler@_V1@sycl@@AEBA?AW4CGType@detail@23@XZ ?getValueFromDynamicParameter@detail@_V1@sycl@@YAPEAXAEAVdynamic_parameter_base@1experimental@oneapi@ext@23@@Z ?get_access_mode@experimental@oneapi@ext@_V1@sycl@@YA?AW4address_access_mode@12345@PEBX_KAEBVcontext@45@@Z -?get_active_index@dynamic_command_group@experimental@oneapi@ext@_V1@sycl@@QEBA_KXZ ?get_addressing_mode@sampler@_V1@sycl@@QEBA?AW4addressing_mode@23@XZ ?get_allocator_internal@buffer_plain@detail@_V1@sycl@@IEBAAEBV?$unique_ptr@VSYCLMemObjAllocator@detail@_V1@sycl@@U?$default_delete@VSYCLMemObjAllocator@detail@_V1@sycl@@@std@@@std@@XZ ?get_allocator_internal@image_plain@detail@_V1@sycl@@IEBAAEBV?$unique_ptr@VSYCLMemObjAllocator@detail@_V1@sycl@@U?$default_delete@VSYCLMemObjAllocator@detail@_V1@sycl@@@std@@@std@@XZ @@ -4275,7 +4277,6 @@ ?setType@handler@_V1@sycl@@AEAAXW4CGType@detail@23@@Z ?setUserFacingNodeType@handler@_V1@sycl@@AEAAXW4node_type@experimental@oneapi@ext@23@@Z ?set_access_mode@experimental@oneapi@ext@_V1@sycl@@YAXPEBX_KW4address_access_mode@12345@AEBVcontext@45@@Z -?set_active_index@dynamic_command_group@experimental@oneapi@ext@_V1@sycl@@QEAAX_K@Z ?set_arg@handler@_V1@sycl@@QEAAXH$$QEAVraw_kernel_arg@experimental@oneapi@ext@23@@Z ?set_final_data_internal@buffer_plain@detail@_V1@sycl@@IEAAXAEBV?$function@$$A6AXAEBV?$function@$$A6AXPEAX@Z@std@@@Z@std@@@Z ?set_final_data_internal@buffer_plain@detail@_V1@sycl@@IEAAXXZ @@ -4307,7 +4308,9 @@ ?submit_impl_and_postprocess@queue@_V1@sycl@@AEAA?AVevent@23@V?$function@$$A6AXAEAVhandler@_V1@sycl@@@Z@std@@AEBUcode_location@detail@23@AEBV?$function@$$A6AX_N0AEAVevent@_V1@sycl@@@Z@6@_N@Z ?submit_impl_and_postprocess@queue@_V1@sycl@@AEAA?AVevent@23@V?$function@$$A6AXAEAVhandler@_V1@sycl@@@Z@std@@V123@AEBUcode_location@detail@23@AEBV?$function@$$A6AX_N0AEAVevent@_V1@sycl@@@Z@6@@Z ?submit_impl_and_postprocess@queue@_V1@sycl@@AEAA?AVevent@23@V?$function@$$A6AXAEAVhandler@_V1@sycl@@@Z@std@@V123@AEBUcode_location@detail@23@AEBV?$function@$$A6AX_N0AEAVevent@_V1@sycl@@@Z@6@_N@Z +?submit_with_event_impl@queue@_V1@sycl@@AEAA?AVevent@23@AEBVtype_erased_cgfo_ty@detail@23@AEBVSubmissionInfo@623@AEBUcode_location@623@_N@Z ?submit_with_event_impl@queue@_V1@sycl@@AEAA?AVevent@23@V?$function@$$A6AXAEAVhandler@_V1@sycl@@@Z@std@@AEBVSubmissionInfo@detail@23@AEBUcode_location@823@_N@Z +?submit_without_event_impl@queue@_V1@sycl@@AEAAXAEBVtype_erased_cgfo_ty@detail@23@AEBVSubmissionInfo@523@AEBUcode_location@523@_N@Z ?submit_without_event_impl@queue@_V1@sycl@@AEAAXV?$function@$$A6AXAEAVhandler@_V1@sycl@@@Z@std@@AEBUcode_location@detail@23@@Z ?submit_without_event_impl@queue@_V1@sycl@@AEAAXV?$function@$$A6AXAEAVhandler@_V1@sycl@@@Z@std@@AEBUcode_location@detail@23@_N@Z ?submit_without_event_impl@queue@_V1@sycl@@AEAAXV?$function@$$A6AXAEAVhandler@_V1@sycl@@@Z@std@@AEBVSubmissionInfo@detail@23@AEBUcode_location@723@_N@Z From cc5189e43f22596c6b805b15def47721d77c151b Mon Sep 17 00:00:00 2001 From: Lorenc Bushi Date: Wed, 22 Jan 2025 19:47:31 -0800 Subject: [PATCH 25/53] Apply suggestions --- sycl/include/sycl/kernel_bundle.hpp | 31 +++++++++---------- sycl/source/kernel_bundle.cpp | 16 +++++----- .../DeviceImageBackendContent/basic_test.cpp | 7 ++--- sycl/test/abi/sycl_symbols_linux.dump | 6 ++-- sycl/test/abi/sycl_symbols_windows.dump | 6 ++-- 5 files changed, 31 insertions(+), 35 deletions(-) diff --git a/sycl/include/sycl/kernel_bundle.hpp b/sycl/include/sycl/kernel_bundle.hpp index fc1805a65e290..b60c197dc96eb 100644 --- a/sycl/include/sycl/kernel_bundle.hpp +++ b/sycl/include/sycl/kernel_bundle.hpp @@ -117,16 +117,6 @@ class __SYCL_EXPORT device_image_plain { ur_native_handle_t getNative() const; - backend ext_oneapi_get_backend() const noexcept; - -#if (!defined(_HAS_STD_BYTE) || _HAS_STD_BYTE != 0) - std::vector ext_oneapi_get_backend_content() const; - - std::pair - ext_oneapi_get_backend_content_view() const; - -#endif // HAS_STD_BYTE - protected: detail::DeviceImageImplPtr impl; @@ -136,6 +126,15 @@ class __SYCL_EXPORT device_image_plain { template friend T detail::createSyclObjFromImpl(decltype(T::impl) ImplObj); + + backend ext_oneapi_get_backend_impl() const noexcept; + +#if (!defined(_HAS_STD_BYTE) || _HAS_STD_BYTE != 0) + std::vector ext_oneapi_get_backend_content_impl() const; + + std::pair + ext_oneapi_get_backend_content_view_impl() const; +#endif // HAS_STD_BYTE }; } // namespace detail @@ -159,23 +158,23 @@ class device_image : public detail::device_image_plain, } backend ext_oneapi_get_backend() const noexcept { - return device_image_plain::ext_oneapi_get_backend(); + return device_image_plain::ext_oneapi_get_backend_impl(); } #if (!defined(_HAS_STD_BYTE) || _HAS_STD_BYTE != 0) template > std::vector ext_oneapi_get_backend_content() const { - return device_image_plain::ext_oneapi_get_backend_content(); + return device_image_plain::ext_oneapi_get_backend_content_impl(); } #ifdef __cpp_lib_span template > - std::span ext_oneapi_get_content_backend_view() const { - return std::span{ - device_image_plain::ext_oneapi_get_backend_content_view().first, - device_image_plain::ext_oneapi_get_backend_content_view().second}; + std::span ext_oneapi_get_backend_content_view() const { + const auto view = + device_image_plain::ext_oneapi_get_backend_content_view_impl(); + return std::span{view.first, view.second}; } #endif // __cpp_lib_span #endif // _HAS_STD_BYTE diff --git a/sycl/source/kernel_bundle.cpp b/sycl/source/kernel_bundle.cpp index bdaea463cf3cb..ff3f23c3bd181 100644 --- a/sycl/source/kernel_bundle.cpp +++ b/sycl/source/kernel_bundle.cpp @@ -45,25 +45,25 @@ ur_native_handle_t device_image_plain::getNative() const { return impl->getNative(); } -backend device_image_plain::ext_oneapi_get_backend() const noexcept { +backend device_image_plain::ext_oneapi_get_backend_impl() const noexcept { return impl->get_context().get_backend(); } std::vector -device_image_plain::ext_oneapi_get_backend_content() const { +device_image_plain::ext_oneapi_get_backend_content_impl() const { return std::vector(reinterpret_cast( impl->get_bin_image_ref()->getRawData().BinaryStart), reinterpret_cast( impl->get_bin_image_ref()->getRawData().BinaryEnd)); } -std::pair -device_image_plain::ext_oneapi_get_backend_content_view() const { +std::pair +device_image_plain::ext_oneapi_get_backend_content_view_impl() const { return std::make_pair( - reinterpret_cast( - impl->get_bin_image_ref()->getRawData().BinaryStart), - reinterpret_cast( - impl->get_bin_image_ref()->getRawData().BinaryEnd)); + const_cast(reinterpret_cast( + impl->get_bin_image_ref()->getRawData().BinaryStart)), + const_cast(reinterpret_cast( + impl->get_bin_image_ref()->getRawData().BinaryEnd))); } //////////////////////////// diff --git a/sycl/test-e2e/DeviceImageBackendContent/basic_test.cpp b/sycl/test-e2e/DeviceImageBackendContent/basic_test.cpp index da87d3d1edff9..395697a87e762 100644 --- a/sycl/test-e2e/DeviceImageBackendContent/basic_test.cpp +++ b/sycl/test-e2e/DeviceImageBackendContent/basic_test.cpp @@ -1,9 +1,6 @@ -// RUN: %{build} -o %t.out +// RUN: %clangxx -fsycl -std=c++20 %s -o %t.out // RUN: %{run} %t.out -#if __has_include() -#include -#endif #include #include #include @@ -48,7 +45,7 @@ int main() { bytes = img.ext_oneapi_get_backend_content(); #ifdef __cpp_lib_span static_assert( - std ::is_same_v); bytes_view = img.ext_oneapi_get_backend_content_view(); assert(bytes_view.size() == bytes.size()); diff --git a/sycl/test/abi/sycl_symbols_linux.dump b/sycl/test/abi/sycl_symbols_linux.dump index b295669c0ac91..b947f3e1ae7fc 100644 --- a/sycl/test/abi/sycl_symbols_linux.dump +++ b/sycl/test/abi/sycl_symbols_linux.dump @@ -3739,9 +3739,9 @@ _ZNK4sycl3_V16detail16AccessorBaseHost6getPtrEv _ZNK4sycl3_V16detail16AccessorBaseHost9getOffsetEv _ZNK4sycl3_V16detail18device_image_plain10has_kernelERKNS0_9kernel_idE _ZNK4sycl3_V16detail18device_image_plain10has_kernelERKNS0_9kernel_idERKNS0_6deviceE -_ZNK4sycl3_V16detail18device_image_plain22ext_oneapi_get_backendEv -_ZNK4sycl3_V16detail18device_image_plain30ext_oneapi_get_backend_contentEv -_ZNK4sycl3_V16detail18device_image_plain35ext_oneapi_get_backend_content_viewEv +_ZNK4sycl3_V16detail18device_image_plain22ext_oneapi_get_backend_implEv +_ZNK4sycl3_V16detail18device_image_plain30ext_oneapi_get_backend_content_implEv +_ZNK4sycl3_V16detail18device_image_plain35ext_oneapi_get_backend_content_view_implEv _ZNK4sycl3_V16detail18device_image_plain9getNativeEv _ZNK4sycl3_V16detail19kernel_bundle_plain10get_kernelERKNS0_9kernel_idE _ZNK4sycl3_V16detail19kernel_bundle_plain10has_kernelERKNS0_9kernel_idE diff --git a/sycl/test/abi/sycl_symbols_windows.dump b/sycl/test/abi/sycl_symbols_windows.dump index a233b2214670d..ccffa70f3560a 100644 --- a/sycl/test/abi/sycl_symbols_windows.dump +++ b/sycl/test/abi/sycl_symbols_windows.dump @@ -3842,9 +3842,9 @@ ?ext_oneapi_empty@queue@_V1@sycl@@QEBA_NXZ ?ext_oneapi_enable_peer_access@device@_V1@sycl@@QEAAXAEBV123@@Z ?ext_oneapi_fill2d_impl@handler@_V1@sycl@@AEAAXPEAX_KPEBX111@Z -?ext_oneapi_get_backend@device_image_plain@detail@_V1@sycl@@QEBA?AW4backend@34@XZ -?ext_oneapi_get_backend_content@device_image_plain@detail@_V1@sycl@@QEBA?AV?$vector@W4byte@std@@V?$allocator@W4byte@std@@@2@@std@@XZ -?ext_oneapi_get_backend_content_view@device_image_plain@detail@_V1@sycl@@QEBA?AU?$pair@PEBW4byte@std@@PEBW412@@std@@XZ +?ext_oneapi_get_backend_impl@device_image_plain@detail@_V1@sycl@@QEBA?AW4backend@34@XZ +?ext_oneapi_get_backend_content_impl@device_image_plain@detail@_V1@sycl@@QEBA?AV?$vector@W4byte@std@@V?$allocator@W4byte@std@@@2@@std@@XZ +?ext_oneapi_get_backend_content_view_impl@device_image_plain@detail@_V1@sycl@@QEBA?AU?$pair@PEBW4byte@std@@PEBW412@@std@@XZ ?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 ?ext_oneapi_get_graph@queue@_V1@sycl@@QEBA?AV?$command_graph@$0A@@experimental@oneapi@ext@23@XZ From 7b8458cc6a7df905d577cabc9f0a7a27378fdd85 Mon Sep 17 00:00:00 2001 From: Lorenc Bushi Date: Wed, 22 Jan 2025 20:37:09 -0800 Subject: [PATCH 26/53] Fix symbols --- sycl/test/abi/sycl_symbols_linux.dump | 2 +- sycl/test/abi/sycl_symbols_windows.dump | 2 +- 2 files changed, 2 insertions(+), 2 deletions(-) diff --git a/sycl/test/abi/sycl_symbols_linux.dump b/sycl/test/abi/sycl_symbols_linux.dump index b947f3e1ae7fc..754cd64373732 100644 --- a/sycl/test/abi/sycl_symbols_linux.dump +++ b/sycl/test/abi/sycl_symbols_linux.dump @@ -3739,9 +3739,9 @@ _ZNK4sycl3_V16detail16AccessorBaseHost6getPtrEv _ZNK4sycl3_V16detail16AccessorBaseHost9getOffsetEv _ZNK4sycl3_V16detail18device_image_plain10has_kernelERKNS0_9kernel_idE _ZNK4sycl3_V16detail18device_image_plain10has_kernelERKNS0_9kernel_idERKNS0_6deviceE -_ZNK4sycl3_V16detail18device_image_plain22ext_oneapi_get_backend_implEv _ZNK4sycl3_V16detail18device_image_plain30ext_oneapi_get_backend_content_implEv _ZNK4sycl3_V16detail18device_image_plain35ext_oneapi_get_backend_content_view_implEv +_ZNK4sycl3_V16detail18device_image_plain22ext_oneapi_get_backend_implEv _ZNK4sycl3_V16detail18device_image_plain9getNativeEv _ZNK4sycl3_V16detail19kernel_bundle_plain10get_kernelERKNS0_9kernel_idE _ZNK4sycl3_V16detail19kernel_bundle_plain10has_kernelERKNS0_9kernel_idE diff --git a/sycl/test/abi/sycl_symbols_windows.dump b/sycl/test/abi/sycl_symbols_windows.dump index ccffa70f3560a..7d31819d7303c 100644 --- a/sycl/test/abi/sycl_symbols_windows.dump +++ b/sycl/test/abi/sycl_symbols_windows.dump @@ -3842,8 +3842,8 @@ ?ext_oneapi_empty@queue@_V1@sycl@@QEBA_NXZ ?ext_oneapi_enable_peer_access@device@_V1@sycl@@QEAAXAEBV123@@Z ?ext_oneapi_fill2d_impl@handler@_V1@sycl@@AEAAXPEAX_KPEBX111@Z -?ext_oneapi_get_backend_impl@device_image_plain@detail@_V1@sycl@@QEBA?AW4backend@34@XZ ?ext_oneapi_get_backend_content_impl@device_image_plain@detail@_V1@sycl@@QEBA?AV?$vector@W4byte@std@@V?$allocator@W4byte@std@@@2@@std@@XZ +?ext_oneapi_get_backend_impl@device_image_plain@detail@_V1@sycl@@QEBA?AW4backend@34@XZ ?ext_oneapi_get_backend_content_view_impl@device_image_plain@detail@_V1@sycl@@QEBA?AU?$pair@PEBW4byte@std@@PEBW412@@std@@XZ ?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 From ab163f735604146f1d125693039ec3371071cddf Mon Sep 17 00:00:00 2001 From: Lorenc Bushi Date: Wed, 22 Jan 2025 20:50:19 -0800 Subject: [PATCH 27/53] Add c++20 flag to compilation of tests --- .../test/extensions/DeviceImageBackendContent/negative_test.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/sycl/test/extensions/DeviceImageBackendContent/negative_test.cpp b/sycl/test/extensions/DeviceImageBackendContent/negative_test.cpp index e48031f7767ee..acfe58053727c 100644 --- a/sycl/test/extensions/DeviceImageBackendContent/negative_test.cpp +++ b/sycl/test/extensions/DeviceImageBackendContent/negative_test.cpp @@ -1,4 +1,4 @@ -// RUN: %clang -fsycl -fsyntax-only -Xclang -verify -Xclang -verify-ignore-unexpected=note,warning %s +// RUN: %clang -fsycl -fsyntax-only -std=c++20 -Xclang -verify -Xclang -verify-ignore-unexpected=note,warning %s #include From 247932e76c71df7bb10120b75cb9983e0a8b5cca Mon Sep 17 00:00:00 2001 From: Lorenc Bushi Date: Wed, 22 Jan 2025 23:50:55 -0500 Subject: [PATCH 28/53] Update negative_test.cpp --- .../test/extensions/DeviceImageBackendContent/negative_test.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/sycl/test/extensions/DeviceImageBackendContent/negative_test.cpp b/sycl/test/extensions/DeviceImageBackendContent/negative_test.cpp index acfe58053727c..bc5783e64ba64 100644 --- a/sycl/test/extensions/DeviceImageBackendContent/negative_test.cpp +++ b/sycl/test/extensions/DeviceImageBackendContent/negative_test.cpp @@ -1,4 +1,4 @@ -// RUN: %clang -fsycl -fsyntax-only -std=c++20 -Xclang -verify -Xclang -verify-ignore-unexpected=note,warning %s +// RUN: %clang -fsycl -fsyntax-only -std=c++20 -Xclang -verify -Xclang -verify-ignore-unexpected=note,warning %s #include From 688521d825e96f1373c560d8c4c8e6a9c56f1c39 Mon Sep 17 00:00:00 2001 From: Lorenc Bushi Date: Thu, 23 Jan 2025 00:16:48 -0500 Subject: [PATCH 29/53] Update sycl_symbols_linux.dump --- sycl/test/abi/sycl_symbols_linux.dump | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/sycl/test/abi/sycl_symbols_linux.dump b/sycl/test/abi/sycl_symbols_linux.dump index 754cd64373732..1d5175482245a 100644 --- a/sycl/test/abi/sycl_symbols_linux.dump +++ b/sycl/test/abi/sycl_symbols_linux.dump @@ -3739,9 +3739,9 @@ _ZNK4sycl3_V16detail16AccessorBaseHost6getPtrEv _ZNK4sycl3_V16detail16AccessorBaseHost9getOffsetEv _ZNK4sycl3_V16detail18device_image_plain10has_kernelERKNS0_9kernel_idE _ZNK4sycl3_V16detail18device_image_plain10has_kernelERKNS0_9kernel_idERKNS0_6deviceE -_ZNK4sycl3_V16detail18device_image_plain30ext_oneapi_get_backend_content_implEv _ZNK4sycl3_V16detail18device_image_plain35ext_oneapi_get_backend_content_view_implEv _ZNK4sycl3_V16detail18device_image_plain22ext_oneapi_get_backend_implEv +_ZNK4sycl3_V16detail18device_image_plain30ext_oneapi_get_backend_content_implEv _ZNK4sycl3_V16detail18device_image_plain9getNativeEv _ZNK4sycl3_V16detail19kernel_bundle_plain10get_kernelERKNS0_9kernel_idE _ZNK4sycl3_V16detail19kernel_bundle_plain10has_kernelERKNS0_9kernel_idE From c1a5ea7e1dec7fe9995b36f92f3eee63f8a0bf93 Mon Sep 17 00:00:00 2001 From: Lorenc Bushi Date: Thu, 23 Jan 2025 00:25:46 -0500 Subject: [PATCH 30/53] Update sycl_symbols_linux.dump --- sycl/test/abi/sycl_symbols_linux.dump | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/sycl/test/abi/sycl_symbols_linux.dump b/sycl/test/abi/sycl_symbols_linux.dump index 1d5175482245a..c27c790da1f6a 100644 --- a/sycl/test/abi/sycl_symbols_linux.dump +++ b/sycl/test/abi/sycl_symbols_linux.dump @@ -3739,9 +3739,9 @@ _ZNK4sycl3_V16detail16AccessorBaseHost6getPtrEv _ZNK4sycl3_V16detail16AccessorBaseHost9getOffsetEv _ZNK4sycl3_V16detail18device_image_plain10has_kernelERKNS0_9kernel_idE _ZNK4sycl3_V16detail18device_image_plain10has_kernelERKNS0_9kernel_idERKNS0_6deviceE -_ZNK4sycl3_V16detail18device_image_plain35ext_oneapi_get_backend_content_view_implEv -_ZNK4sycl3_V16detail18device_image_plain22ext_oneapi_get_backend_implEv -_ZNK4sycl3_V16detail18device_image_plain30ext_oneapi_get_backend_content_implEv +_ZNK4sycl3_V16detail18device_image_plain27ext_oneapi_get_backend_implEv +_ZNK4sycl3_V16detail18device_image_plain40ext_oneapi_get_backend_content_view_implEv +_ZNK4sycl3_V16detail18device_image_plain35ext_oneapi_get_backend_content_implEv _ZNK4sycl3_V16detail18device_image_plain9getNativeEv _ZNK4sycl3_V16detail19kernel_bundle_plain10get_kernelERKNS0_9kernel_idE _ZNK4sycl3_V16detail19kernel_bundle_plain10has_kernelERKNS0_9kernel_idE From bb5b787e147b9066af08081223ba0f609052c7e8 Mon Sep 17 00:00:00 2001 From: Lorenc Bushi Date: Thu, 23 Jan 2025 00:52:44 -0500 Subject: [PATCH 31/53] Update basic_test.cpp --- .../test-e2e/DeviceImageBackendContent/basic_test.cpp | 11 ++++------- 1 file changed, 4 insertions(+), 7 deletions(-) diff --git a/sycl/test-e2e/DeviceImageBackendContent/basic_test.cpp b/sycl/test-e2e/DeviceImageBackendContent/basic_test.cpp index 395697a87e762..6db7dbeabb68e 100644 --- a/sycl/test-e2e/DeviceImageBackendContent/basic_test.cpp +++ b/sycl/test-e2e/DeviceImageBackendContent/basic_test.cpp @@ -7,7 +7,10 @@ class kernel; -void define_kernel(sycl::queue &q) { +int main() { + sycl::device d; + sycl::queue q{d}; + sycl::context ctxt = q.get_context(); int data; sycl::buffer data_buf(&data, 1); q.submit([&](sycl::handler &cgh) { @@ -16,12 +19,6 @@ void define_kernel(sycl::queue &q) { sycl::nd_range{{1}, {1}}, [=](sycl::nd_item<> it) { data_acc[0] = 42; }); }); -} - -int main() { - sycl::device d; - sycl::queue q{d}; - sycl::context ctxt = q.get_context(); sycl::kernel_id id = sycl::get_kernel_id(); auto bundle = sycl::get_kernel_bundle(ctxt, {id}); From a3f837d128bc1c73588a5c68aa63cdcb2cf19149 Mon Sep 17 00:00:00 2001 From: Lorenc Bushi Date: Thu, 23 Jan 2025 01:19:56 -0500 Subject: [PATCH 32/53] Update basic_test.cpp --- sycl/test-e2e/DeviceImageBackendContent/basic_test.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/sycl/test-e2e/DeviceImageBackendContent/basic_test.cpp b/sycl/test-e2e/DeviceImageBackendContent/basic_test.cpp index 6db7dbeabb68e..d02445f65bf33 100644 --- a/sycl/test-e2e/DeviceImageBackendContent/basic_test.cpp +++ b/sycl/test-e2e/DeviceImageBackendContent/basic_test.cpp @@ -21,7 +21,7 @@ int main() { }); sycl::kernel_id id = sycl::get_kernel_id(); auto bundle = - sycl::get_kernel_bundle(ctxt, {id}); + sycl::get_kernel_bundle(ctxt, {d}, {id}); assert(!bundle.empty()); sycl::backend backend; std::vector bytes; From d6a2ca54ab6d0e5a9bce86eced31a5e18606f6c1 Mon Sep 17 00:00:00 2001 From: Lorenc Bushi Date: Thu, 23 Jan 2025 07:20:02 -0800 Subject: [PATCH 33/53] Modify spec and add windows symbols --- .../sycl_ext_oneapi_device_image_backend_content.asciidoc | 6 +++--- sycl/test/abi/sycl_symbols_windows.dump | 2 +- 2 files changed, 4 insertions(+), 4 deletions(-) diff --git a/sycl/doc/extensions/experimental/sycl_ext_oneapi_device_image_backend_content.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_oneapi_device_image_backend_content.asciidoc index 2f581e220873a..c9be4428d25b0 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_device_image_backend_content.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_device_image_backend_content.asciidoc @@ -100,7 +100,7 @@ class device_image { std::vector ext_oneapi_get_backend_content() const; #if defined(__cpp_lib_span) - std::span ext_oneapi_get_backend_content_view() const; + std::span ext_oneapi_get_backend_content_view() const; #endif /*...*/ @@ -148,7 +148,7 @@ See below for a description of the formats used by {dpcpp}. a@ [source,c++] ---- -std::span ext_oneapi_get_content_backend_view() const; +std::span ext_oneapi_get_content_backend_view() const; ---- !==== @@ -157,7 +157,7 @@ Available only when the compiler defines the `__cpp_lib_span` feature-test macro _Constraints:_ Available only when `State` is `bundle_state::executable`. -_Returns:_ A view of the raw backend content for this device image. +_Returns:_ An immutable view of the raw backend content for this device image. The data behind this view has the same lifetime as the `device_image` object. The format of this data is implementation-defined. See below for a description of the formats used by {dpcpp}. diff --git a/sycl/test/abi/sycl_symbols_windows.dump b/sycl/test/abi/sycl_symbols_windows.dump index 7d31819d7303c..ccffa70f3560a 100644 --- a/sycl/test/abi/sycl_symbols_windows.dump +++ b/sycl/test/abi/sycl_symbols_windows.dump @@ -3842,8 +3842,8 @@ ?ext_oneapi_empty@queue@_V1@sycl@@QEBA_NXZ ?ext_oneapi_enable_peer_access@device@_V1@sycl@@QEAAXAEBV123@@Z ?ext_oneapi_fill2d_impl@handler@_V1@sycl@@AEAAXPEAX_KPEBX111@Z -?ext_oneapi_get_backend_content_impl@device_image_plain@detail@_V1@sycl@@QEBA?AV?$vector@W4byte@std@@V?$allocator@W4byte@std@@@2@@std@@XZ ?ext_oneapi_get_backend_impl@device_image_plain@detail@_V1@sycl@@QEBA?AW4backend@34@XZ +?ext_oneapi_get_backend_content_impl@device_image_plain@detail@_V1@sycl@@QEBA?AV?$vector@W4byte@std@@V?$allocator@W4byte@std@@@2@@std@@XZ ?ext_oneapi_get_backend_content_view_impl@device_image_plain@detail@_V1@sycl@@QEBA?AU?$pair@PEBW4byte@std@@PEBW412@@std@@XZ ?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 From d5c214c01d42958d092b0b7bbeacaded2aafe7c0 Mon Sep 17 00:00:00 2001 From: Lorenc Bushi Date: Thu, 23 Jan 2025 08:11:13 -0800 Subject: [PATCH 34/53] Add windows symbols --- sycl/test/abi/sycl_symbols_windows.dump | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/sycl/test/abi/sycl_symbols_windows.dump b/sycl/test/abi/sycl_symbols_windows.dump index ccffa70f3560a..0d88e9b8ccf8b 100644 --- a/sycl/test/abi/sycl_symbols_windows.dump +++ b/sycl/test/abi/sycl_symbols_windows.dump @@ -3842,9 +3842,9 @@ ?ext_oneapi_empty@queue@_V1@sycl@@QEBA_NXZ ?ext_oneapi_enable_peer_access@device@_V1@sycl@@QEAAXAEBV123@@Z ?ext_oneapi_fill2d_impl@handler@_V1@sycl@@AEAAXPEAX_KPEBX111@Z -?ext_oneapi_get_backend_impl@device_image_plain@detail@_V1@sycl@@QEBA?AW4backend@34@XZ -?ext_oneapi_get_backend_content_impl@device_image_plain@detail@_V1@sycl@@QEBA?AV?$vector@W4byte@std@@V?$allocator@W4byte@std@@@2@@std@@XZ -?ext_oneapi_get_backend_content_view_impl@device_image_plain@detail@_V1@sycl@@QEBA?AU?$pair@PEBW4byte@std@@PEBW412@@std@@XZ +?ext_oneapi_get_backend_impl@device_image_plain@detail@_V1@sycl@@IEBA?AW4backend@34@XZ +?ext_oneapi_get_backend_content_impl@device_image_plain@detail@_V1@sycl@@IEBA?AV?$vector@W4byte@std@@V?$allocator@W4byte@std@@@2@@std@@XZ +?ext_oneapi_get_backend_content_view_impl@device_image_plain@detail@_V1@sycl@@IEBA?AU?$pair@PEAW4byte@std@@PEAW412@@std@@XZ ?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 ?ext_oneapi_get_graph@queue@_V1@sycl@@QEBA?AV?$command_graph@$0A@@experimental@oneapi@ext@23@XZ From 2e3811772dc2884cadfe82b1de4e625bc7ae6a06 Mon Sep 17 00:00:00 2001 From: Lorenc Bushi Date: Thu, 23 Jan 2025 13:53:24 -0500 Subject: [PATCH 35/53] Update basic_test.cpp --- sycl/test-e2e/DeviceImageBackendContent/basic_test.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/sycl/test-e2e/DeviceImageBackendContent/basic_test.cpp b/sycl/test-e2e/DeviceImageBackendContent/basic_test.cpp index d02445f65bf33..53fa283d941ed 100644 --- a/sycl/test-e2e/DeviceImageBackendContent/basic_test.cpp +++ b/sycl/test-e2e/DeviceImageBackendContent/basic_test.cpp @@ -1,4 +1,4 @@ -// RUN: %clangxx -fsycl -std=c++20 %s -o %t.out +// RUN: %clangxx -fsycl %s -o %t.out // RUN: %{run} %t.out #include From 4ee145e3d98ba96ef31b594ed48f9c7ca28ae3c2 Mon Sep 17 00:00:00 2001 From: Lorenc Bushi Date: Thu, 23 Jan 2025 14:11:27 -0500 Subject: [PATCH 36/53] Update basic_test.cpp --- sycl/test-e2e/DeviceImageBackendContent/basic_test.cpp | 4 +--- 1 file changed, 1 insertion(+), 3 deletions(-) diff --git a/sycl/test-e2e/DeviceImageBackendContent/basic_test.cpp b/sycl/test-e2e/DeviceImageBackendContent/basic_test.cpp index 53fa283d941ed..2fbf8ff49e48b 100644 --- a/sycl/test-e2e/DeviceImageBackendContent/basic_test.cpp +++ b/sycl/test-e2e/DeviceImageBackendContent/basic_test.cpp @@ -1,12 +1,10 @@ -// RUN: %clangxx -fsycl %s -o %t.out +// RUN: %clangxx -fsycl -std=c++20 %s -o %t.out // RUN: %{run} %t.out #include #include #include -class kernel; - int main() { sycl::device d; sycl::queue q{d}; From 3308fac869c4f48418d7a89d048a0298c3a17d34 Mon Sep 17 00:00:00 2001 From: Lorenc Bushi Date: Thu, 23 Jan 2025 11:29:05 -0800 Subject: [PATCH 37/53] Simplify basic_test.cpp --- sycl/test-e2e/DeviceImageBackendContent/basic_test.cpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/sycl/test-e2e/DeviceImageBackendContent/basic_test.cpp b/sycl/test-e2e/DeviceImageBackendContent/basic_test.cpp index 2fbf8ff49e48b..9825fa314d454 100644 --- a/sycl/test-e2e/DeviceImageBackendContent/basic_test.cpp +++ b/sycl/test-e2e/DeviceImageBackendContent/basic_test.cpp @@ -17,14 +17,14 @@ int main() { sycl::nd_range{{1}, {1}}, [=](sycl::nd_item<> it) { data_acc[0] = 42; }); }); - sycl::kernel_id id = sycl::get_kernel_id(); + sycl::kernel_id id = sycl::get_kernel_id(); auto bundle = sycl::get_kernel_bundle(ctxt, {d}, {id}); assert(!bundle.empty()); sycl::backend backend; std::vector bytes; #ifdef __cpp_lib_span - std::span bytes_view; + std::span bytes_view; #endif for (const auto &img : bundle) { // Check that all 3 functions of the api return correct types and compile. From bc695cfa7e93b9e06586802cfeec72c12b995e9f Mon Sep 17 00:00:00 2001 From: Lorenc Bushi Date: Thu, 23 Jan 2025 15:01:46 -0500 Subject: [PATCH 38/53] Update sycl_ext_oneapi_device_image_backend_content.asciidoc --- .../sycl_ext_oneapi_device_image_backend_content.asciidoc | 4 +--- 1 file changed, 1 insertion(+), 3 deletions(-) diff --git a/sycl/doc/extensions/experimental/sycl_ext_oneapi_device_image_backend_content.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_oneapi_device_image_backend_content.asciidoc index 5f15c6aaa708a..f811d930e2c26 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_device_image_backend_content.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_device_image_backend_content.asciidoc @@ -99,9 +99,7 @@ class device_image { backend ext_oneapi_get_backend() const noexcept; std::vector ext_oneapi_get_backend_content() const; -#if defined(__cpp_lib_span) - std::span ext_oneapi_get_backend_content_view() const; -#endif + std::span ext_oneapi_get_backend_content_view() const; // Requires C++20 /*...*/ }; From 3b2691f82681a95eb472289103c56c7e06484057 Mon Sep 17 00:00:00 2001 From: Lorenc Bushi Date: Thu, 23 Jan 2025 13:00:07 -0800 Subject: [PATCH 39/53] Make span const --- sycl/include/sycl/kernel_bundle.hpp | 6 +++--- sycl/source/kernel_bundle.cpp | 10 +++++----- 2 files changed, 8 insertions(+), 8 deletions(-) diff --git a/sycl/include/sycl/kernel_bundle.hpp b/sycl/include/sycl/kernel_bundle.hpp index b60c197dc96eb..c58ebd7bd001e 100644 --- a/sycl/include/sycl/kernel_bundle.hpp +++ b/sycl/include/sycl/kernel_bundle.hpp @@ -132,7 +132,7 @@ class __SYCL_EXPORT device_image_plain { #if (!defined(_HAS_STD_BYTE) || _HAS_STD_BYTE != 0) std::vector ext_oneapi_get_backend_content_impl() const; - std::pair + std::pair ext_oneapi_get_backend_content_view_impl() const; #endif // HAS_STD_BYTE }; @@ -171,10 +171,10 @@ class device_image : public detail::device_image_plain, #ifdef __cpp_lib_span template > - std::span ext_oneapi_get_backend_content_view() const { + std::span ext_oneapi_get_backend_content_view() const { const auto view = device_image_plain::ext_oneapi_get_backend_content_view_impl(); - return std::span{view.first, view.second}; + return std::span{view.first, view.second}; } #endif // __cpp_lib_span #endif // _HAS_STD_BYTE diff --git a/sycl/source/kernel_bundle.cpp b/sycl/source/kernel_bundle.cpp index ff3f23c3bd181..5c2d75737f2f7 100644 --- a/sycl/source/kernel_bundle.cpp +++ b/sycl/source/kernel_bundle.cpp @@ -57,13 +57,13 @@ device_image_plain::ext_oneapi_get_backend_content_impl() const { impl->get_bin_image_ref()->getRawData().BinaryEnd)); } -std::pair +std::pair device_image_plain::ext_oneapi_get_backend_content_view_impl() const { return std::make_pair( - const_cast(reinterpret_cast( - impl->get_bin_image_ref()->getRawData().BinaryStart)), - const_cast(reinterpret_cast( - impl->get_bin_image_ref()->getRawData().BinaryEnd))); + reinterpret_cast( + impl->get_bin_image_ref()->getRawData().BinaryStart), + reinterpret_cast( + impl->get_bin_image_ref()->getRawData().BinaryEnd)); } //////////////////////////// From 743c9a71fd904440adc5cb5d2323d0066fb34fb4 Mon Sep 17 00:00:00 2001 From: Lorenc Bushi Date: Fri, 24 Jan 2025 09:58:25 -0800 Subject: [PATCH 40/53] Fix test failures on HIP/CUDA --- .../DeviceImageBackendContent/basic_test.cpp | 19 ++++++++----------- 1 file changed, 8 insertions(+), 11 deletions(-) diff --git a/sycl/test-e2e/DeviceImageBackendContent/basic_test.cpp b/sycl/test-e2e/DeviceImageBackendContent/basic_test.cpp index 9825fa314d454..ff9e2cee8c4be 100644 --- a/sycl/test-e2e/DeviceImageBackendContent/basic_test.cpp +++ b/sycl/test-e2e/DeviceImageBackendContent/basic_test.cpp @@ -6,21 +6,18 @@ #include int main() { - sycl::device d; - sycl::queue q{d}; + sycl::queue q; sycl::context ctxt = q.get_context(); - int data; - sycl::buffer data_buf(&data, 1); - q.submit([&](sycl::handler &cgh) { - sycl::accessor data_acc(data_buf, cgh); - cgh.parallel_for( - sycl::nd_range{{1}, {1}}, - [=](sycl::nd_item<> it) { data_acc[0] = 42; }); - }); + sycl::buffer buf(sycl::range<1>(1)); sycl::kernel_id id = sycl::get_kernel_id(); auto bundle = - sycl::get_kernel_bundle(ctxt, {d}, {id}); + sycl::get_kernel_bundle(ctxt, {id}); assert(!bundle.empty()); + sycl::kernel krn = bundle.get_kernel(id); + q.submit([&](sycl::handler &cgh) { + sycl::accessor acc(buf, cgh); + cgh.single_task(krn, [=]() { acc[0] = 42; }); + }); sycl::backend backend; std::vector bytes; #ifdef __cpp_lib_span From 97637cfd99b2249fc21596e295d9fedb1efe803f Mon Sep 17 00:00:00 2001 From: Lorenc Bushi Date: Fri, 24 Jan 2025 10:46:56 -0800 Subject: [PATCH 41/53] Fix test failures on HIP/CUDA --- sycl/test-e2e/DeviceImageBackendContent/basic_test.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/sycl/test-e2e/DeviceImageBackendContent/basic_test.cpp b/sycl/test-e2e/DeviceImageBackendContent/basic_test.cpp index ff9e2cee8c4be..f1e84273604eb 100644 --- a/sycl/test-e2e/DeviceImageBackendContent/basic_test.cpp +++ b/sycl/test-e2e/DeviceImageBackendContent/basic_test.cpp @@ -1,4 +1,4 @@ -// RUN: %clangxx -fsycl -std=c++20 %s -o %t.out +// RUN: %{build} -std=c++20 -o %t.out // RUN: %{run} %t.out #include From afe72b866d1e9c14ecc91f4e7803dadb2064e93c Mon Sep 17 00:00:00 2001 From: Lorenc Bushi Date: Fri, 24 Jan 2025 10:49:23 -0800 Subject: [PATCH 42/53] Fix test failures on HIP/CUDA --- sycl/test-e2e/DeviceImageBackendContent/basic_test.cpp | 8 ++++---- 1 file changed, 4 insertions(+), 4 deletions(-) diff --git a/sycl/test-e2e/DeviceImageBackendContent/basic_test.cpp b/sycl/test-e2e/DeviceImageBackendContent/basic_test.cpp index f1e84273604eb..8d0c873f1d2d3 100644 --- a/sycl/test-e2e/DeviceImageBackendContent/basic_test.cpp +++ b/sycl/test-e2e/DeviceImageBackendContent/basic_test.cpp @@ -9,14 +9,14 @@ int main() { sycl::queue q; sycl::context ctxt = q.get_context(); sycl::buffer buf(sycl::range<1>(1)); - sycl::kernel_id id = sycl::get_kernel_id(); + sycl::kernel_id k_id = sycl::get_kernel_id(); auto bundle = - sycl::get_kernel_bundle(ctxt, {id}); + sycl::get_kernel_bundle(ctxt, {k_id}); assert(!bundle.empty()); - sycl::kernel krn = bundle.get_kernel(id); + sycl::kernel krn = bundle.get_kernel(k_id); q.submit([&](sycl::handler &cgh) { sycl::accessor acc(buf, cgh); - cgh.single_task(krn, [=]() { acc[0] = 42; }); + cgh.single_task(krn, [=]() { acc[0] = 42; }); }); sycl::backend backend; std::vector bytes; From d1f26443a64fc1e3f67f7ca145d9f1f03855f8d9 Mon Sep 17 00:00:00 2001 From: Lorenc Bushi Date: Fri, 24 Jan 2025 15:08:09 -0500 Subject: [PATCH 43/53] Update sycl_symbols_windows.dump --- sycl/test/abi/sycl_symbols_windows.dump | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/sycl/test/abi/sycl_symbols_windows.dump b/sycl/test/abi/sycl_symbols_windows.dump index 0d88e9b8ccf8b..88b14196d5f9a 100644 --- a/sycl/test/abi/sycl_symbols_windows.dump +++ b/sycl/test/abi/sycl_symbols_windows.dump @@ -3844,7 +3844,7 @@ ?ext_oneapi_fill2d_impl@handler@_V1@sycl@@AEAAXPEAX_KPEBX111@Z ?ext_oneapi_get_backend_impl@device_image_plain@detail@_V1@sycl@@IEBA?AW4backend@34@XZ ?ext_oneapi_get_backend_content_impl@device_image_plain@detail@_V1@sycl@@IEBA?AV?$vector@W4byte@std@@V?$allocator@W4byte@std@@@2@@std@@XZ -?ext_oneapi_get_backend_content_view_impl@device_image_plain@detail@_V1@sycl@@IEBA?AU?$pair@PEAW4byte@std@@PEAW412@@std@@XZ +?ext_oneapi_get_backend_content_view_impl@device_image_plain@detail@_V1@sycl@@IEBA?AU?$pair@PEBW4byte@std@@PEBW412@@std@@XZ ?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 ?ext_oneapi_get_graph@queue@_V1@sycl@@QEBA?AV?$command_graph@$0A@@experimental@oneapi@ext@23@XZ From 910ce34143c7a8e9c2185497505e6e9f28daa20f Mon Sep 17 00:00:00 2001 From: Lorenc Bushi Date: Tue, 28 Jan 2025 12:03:15 -0800 Subject: [PATCH 44/53] Add OpenCL interoperability test --- .../OCL_interop_test.cpp | 81 +++++++++++++++++++ 1 file changed, 81 insertions(+) create mode 100644 sycl/test-e2e/DeviceImageBackendContent/OCL_interop_test.cpp diff --git a/sycl/test-e2e/DeviceImageBackendContent/OCL_interop_test.cpp b/sycl/test-e2e/DeviceImageBackendContent/OCL_interop_test.cpp new file mode 100644 index 0000000000000..6ef34c4544d29 --- /dev/null +++ b/sycl/test-e2e/DeviceImageBackendContent/OCL_interop_test.cpp @@ -0,0 +1,81 @@ +// REQUIRES: opencl, opencl_icd, aspect-usm_shared_allocations +// RUN: %{build} %opencl_lib -fno-sycl-dead-args-optimization -o %t.out +// RUN: %{run} %t.out +// +#include +#include +#include +#include +#include +#include + +namespace syclext = sycl::ext::oneapi; +namespace syclexp = sycl::ext::oneapi::experimental; + +extern "C" SYCL_EXT_ONEAPI_FUNCTION_PROPERTY( + (syclexp::nd_range_kernel<1>)) void iota(int *ptr) { + size_t id = syclext::this_work_item::get_nd_item<1>().get_global_linear_id(); + ptr[id] = 42; +} + +int main() { + sycl::queue q; + sycl::context ctxt = q.get_context(); + sycl::device d = ctxt.get_devices()[0]; + + // The following ifndef is required due to a number of limitations of free + // function kernels. See CMPLRLLVM-61498. + // TODO: Remove it once these limitations are no longer there. +#ifndef __SYCL_DEVICE_ONLY__ + // First, run the kernel using the SYCL API. + auto Bundle = sycl::get_kernel_bundle(ctxt); + sycl::kernel_id iota_id = syclexp::get_kernel_id(); + sycl::kernel k_iota = Bundle.get_kernel(iota_id); + + int *ptr = sycl::malloc_shared(1, q); + *ptr = 0; + q.submit([&](sycl::handler &cgh) { + cgh.set_args(ptr); + cgh.parallel_for(sycl::nd_range{{1}, {1}}, k_iota); + }).wait(); + // Now, run the kernel by first getting its image as an executable, + // making an OCL kernel out of it and then making a SYCL kernel out of + // the OCL kernel. Run this kernel on the SYCL API and verify + // that it has the same result as the kernel that was run directly on SYCL + // API. First, get a kernel bundle that contains the kernel "iota". + auto exe_bndl = sycl::get_kernel_bundle( + ctxt, {d}, + [&](const sycl::device_image &img) { + return img.has_kernel(iota_id, d); + }); + assert(!exe_bndl.empty()); + std::vector bytes; + const sycl::device_image &img = + *(exe_bndl.begin()); + bytes = img.ext_oneapi_get_backend_content(); + + auto clContext = sycl::get_native(ctxt); + auto clDevice = sycl::get_native(d); + + cl_int status; + auto clProgram = clCreateProgramWithIL( + clContext, reinterpret_cast(bytes.data()), bytes.size(), + &status); + assert(status == CL_SUCCESS); + status = clBuildProgram(clProgram, 1, &clDevice, "", nullptr, nullptr); + assert(status == CL_SUCCESS); + auto clKernel = clCreateKernel(clProgram, "__sycl_kernel_iota", &status); + assert(status == CL_SUCCESS); + sycl::kernel k_iota_twin = + sycl::make_kernel(clKernel, ctxt); + int *ptr_twin = sycl::malloc_shared(1, q); + *ptr_twin = 1; + q.submit([&](sycl::handler &cgh) { + cgh.set_args(ptr_twin); + cgh.parallel_for(sycl::nd_range{{1}, {1}}, k_iota_twin); + }).wait(); + assert(*ptr_twin == *ptr); + sycl::free(ptr, q); + sycl::free(ptr_twin, q); +#endif +} From d267155215865c5934e9919e131cf0e1d33abbef Mon Sep 17 00:00:00 2001 From: Lorenc Bushi Date: Tue, 28 Jan 2025 12:04:56 -0800 Subject: [PATCH 45/53] Add OpenCL interoperability test --- .../test-e2e/DeviceImageBackendContent/OCL_interop_test.cpp | 6 ++++-- 1 file changed, 4 insertions(+), 2 deletions(-) diff --git a/sycl/test-e2e/DeviceImageBackendContent/OCL_interop_test.cpp b/sycl/test-e2e/DeviceImageBackendContent/OCL_interop_test.cpp index 6ef34c4544d29..d878c6efcd11c 100644 --- a/sycl/test-e2e/DeviceImageBackendContent/OCL_interop_test.cpp +++ b/sycl/test-e2e/DeviceImageBackendContent/OCL_interop_test.cpp @@ -19,9 +19,11 @@ extern "C" SYCL_EXT_ONEAPI_FUNCTION_PROPERTY( } int main() { - sycl::queue q; + sycl::device d([](const sycl::device &d) { + return d.get_backend() == sycl::backend::opencl; + }); + sycl::queue q{d}; sycl::context ctxt = q.get_context(); - sycl::device d = ctxt.get_devices()[0]; // The following ifndef is required due to a number of limitations of free // function kernels. See CMPLRLLVM-61498. From 2c916f6bab345e4261e75eb834fd5d47e0d0eeef Mon Sep 17 00:00:00 2001 From: Lorenc Bushi Date: Fri, 31 Jan 2025 13:26:00 -0800 Subject: [PATCH 46/53] Add CUDA test --- .../CUDA_interop_test.cpp | 33 +++++++++++++++++++ 1 file changed, 33 insertions(+) create mode 100644 sycl/test-e2e/DeviceImageBackendContent/CUDA_interop_test.cpp diff --git a/sycl/test-e2e/DeviceImageBackendContent/CUDA_interop_test.cpp b/sycl/test-e2e/DeviceImageBackendContent/CUDA_interop_test.cpp new file mode 100644 index 0000000000000..6000188af1649 --- /dev/null +++ b/sycl/test-e2e/DeviceImageBackendContent/CUDA_interop_test.cpp @@ -0,0 +1,33 @@ +// RUN: %{build} %cuda_options -o %t.out +// RUN: %{run} %t.out +// REQUIRES: cuda, cuda_dev_kit + +#include +#include + +#include + +int main() { + sycl::device d([](const sycl::device &d) { + return d.get_backend() == sycl::backend::opencl; + }); + sycl::queue q{d}; + sycl::context ctxt = q.get_context(); + sycl::kernel_id k_id = sycl::get_kernel_id(); + auto bundle = + sycl::get_kernel_bundle(ctxt, {k_id}); + assert(!bundle.empty()); + sycl::kernel krn = bundle.get_kernel(k_id); + sycl::buffer buf(sycl::range<1>(1)); + q.submit([&](sycl::handler &cgh) { + sycl::accessor acc(buf, cgh); + cgh.single_task(krn, [=]() { acc[0] = 42; }); + }); + const auto img = *(bundle.begin()); + const auto bytes = img.ext_oneapi_get_backend_content(); + CUmodule m; + CUresult res = + cuModuleLoadData(&m, reinterpret_cast(bytes.data())); + assert(result == CUDA_SUCCESS); + return 0; +} From 77eda189dafd0439ab1ff86062a3aa2aa3ee4fa0 Mon Sep 17 00:00:00 2001 From: Lorenc Bushi Date: Fri, 31 Jan 2025 13:28:12 -0800 Subject: [PATCH 47/53] Add CUDA test --- sycl/test-e2e/DeviceImageBackendContent/CUDA_interop_test.cpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/sycl/test-e2e/DeviceImageBackendContent/CUDA_interop_test.cpp b/sycl/test-e2e/DeviceImageBackendContent/CUDA_interop_test.cpp index 6000188af1649..9ba7c11b6d00c 100644 --- a/sycl/test-e2e/DeviceImageBackendContent/CUDA_interop_test.cpp +++ b/sycl/test-e2e/DeviceImageBackendContent/CUDA_interop_test.cpp @@ -2,10 +2,10 @@ // RUN: %{run} %t.out // REQUIRES: cuda, cuda_dev_kit +#include #include #include - -#include +#include int main() { sycl::device d([](const sycl::device &d) { From a29db256117b987d1fa38c8b4915cdda864d44ca Mon Sep 17 00:00:00 2001 From: Lorenc Bushi Date: Wed, 5 Feb 2025 08:25:20 -0800 Subject: [PATCH 48/53] Fix typo in CUDA test --- sycl/test-e2e/DeviceImageBackendContent/CUDA_interop_test.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/sycl/test-e2e/DeviceImageBackendContent/CUDA_interop_test.cpp b/sycl/test-e2e/DeviceImageBackendContent/CUDA_interop_test.cpp index 9ba7c11b6d00c..cea549af54d0d 100644 --- a/sycl/test-e2e/DeviceImageBackendContent/CUDA_interop_test.cpp +++ b/sycl/test-e2e/DeviceImageBackendContent/CUDA_interop_test.cpp @@ -9,7 +9,7 @@ int main() { sycl::device d([](const sycl::device &d) { - return d.get_backend() == sycl::backend::opencl; + return d.get_backend() == sycl::backend::cuda; }); sycl::queue q{d}; sycl::context ctxt = q.get_context(); From 248fc93dc3472f44b6b3acc96eeaa7cbeaf54e06 Mon Sep 17 00:00:00 2001 From: Lorenc Bushi Date: Wed, 5 Feb 2025 13:35:32 -0500 Subject: [PATCH 49/53] Update CUDA_interop_test.cpp --- sycl/test-e2e/DeviceImageBackendContent/CUDA_interop_test.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/sycl/test-e2e/DeviceImageBackendContent/CUDA_interop_test.cpp b/sycl/test-e2e/DeviceImageBackendContent/CUDA_interop_test.cpp index cea549af54d0d..62fb16e18b202 100644 --- a/sycl/test-e2e/DeviceImageBackendContent/CUDA_interop_test.cpp +++ b/sycl/test-e2e/DeviceImageBackendContent/CUDA_interop_test.cpp @@ -9,7 +9,7 @@ int main() { sycl::device d([](const sycl::device &d) { - return d.get_backend() == sycl::backend::cuda; + return d.get_backend() == sycl::backend::ext_oneapi_cuda; }); sycl::queue q{d}; sycl::context ctxt = q.get_context(); From 4ccce5a6f6cfa2ff828a31e398593a994c280623 Mon Sep 17 00:00:00 2001 From: Lorenc Bushi Date: Wed, 5 Feb 2025 16:34:51 -0500 Subject: [PATCH 50/53] Fix compilation error in CUDA_interop_test.cpp --- sycl/test-e2e/DeviceImageBackendContent/CUDA_interop_test.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/sycl/test-e2e/DeviceImageBackendContent/CUDA_interop_test.cpp b/sycl/test-e2e/DeviceImageBackendContent/CUDA_interop_test.cpp index 62fb16e18b202..f8ee368ed3498 100644 --- a/sycl/test-e2e/DeviceImageBackendContent/CUDA_interop_test.cpp +++ b/sycl/test-e2e/DeviceImageBackendContent/CUDA_interop_test.cpp @@ -26,7 +26,7 @@ int main() { const auto img = *(bundle.begin()); const auto bytes = img.ext_oneapi_get_backend_content(); CUmodule m; - CUresult res = + CUresult result = cuModuleLoadData(&m, reinterpret_cast(bytes.data())); assert(result == CUDA_SUCCESS); return 0; From 2e4d7f83e98bf3a391eb25f15841c5576a6c3d26 Mon Sep 17 00:00:00 2001 From: Lorenc Bushi Date: Thu, 6 Feb 2025 07:27:57 -0800 Subject: [PATCH 51/53] Apply suggestions --- sycl/include/sycl/kernel_bundle.hpp | 6 +++--- sycl/source/kernel_bundle.cpp | 8 -------- .../CUDA_interop_test.cpp | 5 +---- .../DeviceImageBackendContent/L0_interop_test.cpp | 9 +++------ .../DeviceImageBackendContent/OCL_interop_test.cpp | 14 ++++++-------- 5 files changed, 13 insertions(+), 29 deletions(-) diff --git a/sycl/include/sycl/kernel_bundle.hpp b/sycl/include/sycl/kernel_bundle.hpp index c58ebd7bd001e..8da50f05c42a6 100644 --- a/sycl/include/sycl/kernel_bundle.hpp +++ b/sycl/include/sycl/kernel_bundle.hpp @@ -130,8 +130,6 @@ class __SYCL_EXPORT device_image_plain { backend ext_oneapi_get_backend_impl() const noexcept; #if (!defined(_HAS_STD_BYTE) || _HAS_STD_BYTE != 0) - std::vector ext_oneapi_get_backend_content_impl() const; - std::pair ext_oneapi_get_backend_content_view_impl() const; #endif // HAS_STD_BYTE @@ -165,7 +163,9 @@ class device_image : public detail::device_image_plain, template > std::vector ext_oneapi_get_backend_content() const { - return device_image_plain::ext_oneapi_get_backend_content_impl(); + const auto view = + device_image_plain::ext_oneapi_get_backend_content_view_impl(); + return std::vector(view.first, view.second); } #ifdef __cpp_lib_span diff --git a/sycl/source/kernel_bundle.cpp b/sycl/source/kernel_bundle.cpp index 5c2d75737f2f7..06a8d564221ad 100644 --- a/sycl/source/kernel_bundle.cpp +++ b/sycl/source/kernel_bundle.cpp @@ -49,14 +49,6 @@ backend device_image_plain::ext_oneapi_get_backend_impl() const noexcept { return impl->get_context().get_backend(); } -std::vector -device_image_plain::ext_oneapi_get_backend_content_impl() const { - return std::vector(reinterpret_cast( - impl->get_bin_image_ref()->getRawData().BinaryStart), - reinterpret_cast( - impl->get_bin_image_ref()->getRawData().BinaryEnd)); -} - std::pair device_image_plain::ext_oneapi_get_backend_content_view_impl() const { return std::make_pair( diff --git a/sycl/test-e2e/DeviceImageBackendContent/CUDA_interop_test.cpp b/sycl/test-e2e/DeviceImageBackendContent/CUDA_interop_test.cpp index f8ee368ed3498..6225bb11cf141 100644 --- a/sycl/test-e2e/DeviceImageBackendContent/CUDA_interop_test.cpp +++ b/sycl/test-e2e/DeviceImageBackendContent/CUDA_interop_test.cpp @@ -8,10 +8,7 @@ #include int main() { - sycl::device d([](const sycl::device &d) { - return d.get_backend() == sycl::backend::ext_oneapi_cuda; - }); - sycl::queue q{d}; + sycl::queue q; sycl::context ctxt = q.get_context(); sycl::kernel_id k_id = sycl::get_kernel_id(); auto bundle = diff --git a/sycl/test-e2e/DeviceImageBackendContent/L0_interop_test.cpp b/sycl/test-e2e/DeviceImageBackendContent/L0_interop_test.cpp index ab0b920d5626d..486cb81048786 100644 --- a/sycl/test-e2e/DeviceImageBackendContent/L0_interop_test.cpp +++ b/sycl/test-e2e/DeviceImageBackendContent/L0_interop_test.cpp @@ -19,10 +19,7 @@ extern "C" SYCL_EXT_ONEAPI_FUNCTION_PROPERTY( } int main() { - sycl::device d([](const sycl::device &d) { - return d.get_backend() == sycl::backend::ext_oneapi_level_zero; - }); - sycl::queue q{d}; + sycl::queue q; sycl::context ctxt = q.get_context(); // The following ifndef is required due to a number of limitations of free @@ -30,9 +27,9 @@ int main() { // TODO: Remove it once these limitations are no longer there. #ifndef __SYCL_DEVICE_ONLY__ // First, run the kernel using the SYCL API. - auto Bundle = sycl::get_kernel_bundle(ctxt); + auto bundle = sycl::get_kernel_bundle(ctxt); sycl::kernel_id iota_id = syclexp::get_kernel_id(); - sycl::kernel k_iota = Bundle.get_kernel(iota_id); + sycl::kernel k_iota = bundle.get_kernel(iota_id); int *ptr = sycl::malloc_shared(1, q); *ptr = 0; q.submit([&](sycl::handler &cgh) { diff --git a/sycl/test-e2e/DeviceImageBackendContent/OCL_interop_test.cpp b/sycl/test-e2e/DeviceImageBackendContent/OCL_interop_test.cpp index d878c6efcd11c..fd7c6b55fbe19 100644 --- a/sycl/test-e2e/DeviceImageBackendContent/OCL_interop_test.cpp +++ b/sycl/test-e2e/DeviceImageBackendContent/OCL_interop_test.cpp @@ -19,20 +19,18 @@ extern "C" SYCL_EXT_ONEAPI_FUNCTION_PROPERTY( } int main() { - sycl::device d([](const sycl::device &d) { - return d.get_backend() == sycl::backend::opencl; - }); - sycl::queue q{d}; + sycl::queue q; sycl::context ctxt = q.get_context(); - + sycl::device d = ctxt.get_devices()[0]; // The following ifndef is required due to a number of limitations of free // function kernels. See CMPLRLLVM-61498. // TODO: Remove it once these limitations are no longer there. #ifndef __SYCL_DEVICE_ONLY__ // First, run the kernel using the SYCL API. - auto Bundle = sycl::get_kernel_bundle(ctxt); + + auto bundle = sycl::get_kernel_bundle(ctxt); sycl::kernel_id iota_id = syclexp::get_kernel_id(); - sycl::kernel k_iota = Bundle.get_kernel(iota_id); + sycl::kernel k_iota = bundle.get_kernel(iota_id); int *ptr = sycl::malloc_shared(1, q); *ptr = 0; @@ -55,7 +53,7 @@ int main() { const sycl::device_image &img = *(exe_bndl.begin()); bytes = img.ext_oneapi_get_backend_content(); - + std::cout << bytes.size() << std::endl; auto clContext = sycl::get_native(ctxt); auto clDevice = sycl::get_native(d); From 9aa2184658fdd361aa10473222ba38632950c73e Mon Sep 17 00:00:00 2001 From: Lorenc Bushi Date: Thu, 6 Feb 2025 08:46:34 -0800 Subject: [PATCH 52/53] Fix compilation error in L0 test --- sycl/test-e2e/DeviceImageBackendContent/L0_interop_test.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/sycl/test-e2e/DeviceImageBackendContent/L0_interop_test.cpp b/sycl/test-e2e/DeviceImageBackendContent/L0_interop_test.cpp index 486cb81048786..26fb68bca1275 100644 --- a/sycl/test-e2e/DeviceImageBackendContent/L0_interop_test.cpp +++ b/sycl/test-e2e/DeviceImageBackendContent/L0_interop_test.cpp @@ -21,7 +21,7 @@ extern "C" SYCL_EXT_ONEAPI_FUNCTION_PROPERTY( int main() { sycl::queue q; sycl::context ctxt = q.get_context(); - + sycl::device d = ctxt.get_devices()[0]; // The following ifndef is required due to a number of limitations of free // function kernels. See CMPLRLLVM-61498. // TODO: Remove it once these limitations are no longer there. From 6a88f93e9b433a9d4063f872f3fcd0a4b4275bac Mon Sep 17 00:00:00 2001 From: Lorenc Bushi Date: Thu, 6 Feb 2025 08:47:49 -0800 Subject: [PATCH 53/53] Adjust symbols --- sycl/test/abi/sycl_symbols_linux.dump | 1 - sycl/test/abi/sycl_symbols_windows.dump | 1 - 2 files changed, 2 deletions(-) diff --git a/sycl/test/abi/sycl_symbols_linux.dump b/sycl/test/abi/sycl_symbols_linux.dump index 3d4dfc2902861..144e77736de98 100644 --- a/sycl/test/abi/sycl_symbols_linux.dump +++ b/sycl/test/abi/sycl_symbols_linux.dump @@ -3744,7 +3744,6 @@ _ZNK4sycl3_V16detail18device_image_plain10has_kernelERKNS0_9kernel_idE _ZNK4sycl3_V16detail18device_image_plain10has_kernelERKNS0_9kernel_idERKNS0_6deviceE _ZNK4sycl3_V16detail18device_image_plain27ext_oneapi_get_backend_implEv _ZNK4sycl3_V16detail18device_image_plain40ext_oneapi_get_backend_content_view_implEv -_ZNK4sycl3_V16detail18device_image_plain35ext_oneapi_get_backend_content_implEv _ZNK4sycl3_V16detail18device_image_plain9getNativeEv _ZNK4sycl3_V16detail19kernel_bundle_plain10get_kernelERKNS0_9kernel_idE _ZNK4sycl3_V16detail19kernel_bundle_plain10has_kernelERKNS0_9kernel_idE diff --git a/sycl/test/abi/sycl_symbols_windows.dump b/sycl/test/abi/sycl_symbols_windows.dump index 2697bb2590c9f..e601a840b50d1 100644 --- a/sycl/test/abi/sycl_symbols_windows.dump +++ b/sycl/test/abi/sycl_symbols_windows.dump @@ -3863,7 +3863,6 @@ ?ext_oneapi_enable_peer_access@device@_V1@sycl@@QEAAXAEBV123@@Z ?ext_oneapi_fill2d_impl@handler@_V1@sycl@@AEAAXPEAX_KPEBX111@Z ?ext_oneapi_get_backend_impl@device_image_plain@detail@_V1@sycl@@IEBA?AW4backend@34@XZ -?ext_oneapi_get_backend_content_impl@device_image_plain@detail@_V1@sycl@@IEBA?AV?$vector@W4byte@std@@V?$allocator@W4byte@std@@@2@@std@@XZ ?ext_oneapi_get_backend_content_view_impl@device_image_plain@detail@_V1@sycl@@IEBA?AU?$pair@PEBW4byte@std@@PEBW412@@std@@XZ ?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