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 92% 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 index 77eb2703ba386..f811d930e2c26 100644 --- 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 @@ -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.* @@ -101,7 +99,7 @@ class device_image { backend ext_oneapi_get_backend() const noexcept; std::vector ext_oneapi_get_backend_content() const; - std::span ext_oneapi_get_backend_content_view() const; // Requires C++20 + std::span ext_oneapi_get_backend_content_view() const; // Requires C++20 /*...*/ }; @@ -148,7 +146,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; ---- !==== @@ -156,7 +154,7 @@ Minimum C++ Version: {cpp}20 _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/include/sycl/kernel_bundle.hpp b/sycl/include/sycl/kernel_bundle.hpp index a61019efdbf5d..8da50f05c42a6 100644 --- a/sycl/include/sycl/kernel_bundle.hpp +++ b/sycl/include/sycl/kernel_bundle.hpp @@ -26,12 +26,15 @@ #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 +#if __has_include() +#include +#endif #include // for string #include // for enable_if_t, remove_refer... #include // for move @@ -123,6 +126,13 @@ 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::pair + ext_oneapi_get_backend_content_view_impl() const; +#endif // HAS_STD_BYTE }; } // namespace detail @@ -145,6 +155,30 @@ 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::ext_oneapi_get_backend_impl(); + } + +#if (!defined(_HAS_STD_BYTE) || _HAS_STD_BYTE != 0) + template > + std::vector ext_oneapi_get_backend_content() const { + const auto view = + device_image_plain::ext_oneapi_get_backend_content_view_impl(); + return std::vector(view.first, view.second); + } + +#ifdef __cpp_lib_span + template > + 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 + private: device_image(detail::DeviceImageImplPtr Impl) : device_image_plain(std::move(Impl)) {} diff --git a/sycl/source/feature_test.hpp.in b/sycl/source/feature_test.hpp.in index f062302257e4b..7aa9b676626f4 100644 --- a/sycl/source/feature_test.hpp.in +++ b/sycl/source/feature_test.hpp.in @@ -113,6 +113,7 @@ inline namespace _V1 { #define SYCL_EXT_ONEAPI_WORK_GROUP_SCRATCH_MEMORY 1 #define SYCL_EXT_ONEAPI_WORK_GROUP_STATIC 1 #define SYCL_EXT_ONEAPI_NUM_COMPUTE_UNITS 1 +#define SYCL_EXT_ONEAPI_DEVICE_IMAGE_BACKEND_CONTENT 1 // In progress yet #define SYCL_EXT_ONEAPI_ATOMIC16 0 diff --git a/sycl/source/kernel_bundle.cpp b/sycl/source/kernel_bundle.cpp index e19c2b9df2a75..06a8d564221ad 100644 --- a/sycl/source/kernel_bundle.cpp +++ b/sycl/source/kernel_bundle.cpp @@ -45,6 +45,19 @@ ur_native_handle_t device_image_plain::getNative() const { return impl->getNative(); } +backend device_image_plain::ext_oneapi_get_backend_impl() const noexcept { + return impl->get_context().get_backend(); +} + +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)); +} + //////////////////////////// ///// kernel_bundle_plain /////////////////////////// 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..6225bb11cf141 --- /dev/null +++ b/sycl/test-e2e/DeviceImageBackendContent/CUDA_interop_test.cpp @@ -0,0 +1,30 @@ +// RUN: %{build} %cuda_options -o %t.out +// RUN: %{run} %t.out +// REQUIRES: cuda, cuda_dev_kit + +#include +#include +#include +#include + +int main() { + sycl::queue q; + 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 result = + cuModuleLoadData(&m, reinterpret_cast(bytes.data())); + assert(result == CUDA_SUCCESS); + return 0; +} 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..26fb68bca1275 --- /dev/null +++ b/sycl/test-e2e/DeviceImageBackendContent/L0_interop_test.cpp @@ -0,0 +1,94 @@ +// 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::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 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); + }); + assert(!exe_bndl.empty()); + 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); + *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 +} 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..fd7c6b55fbe19 --- /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(); + std::cout << bytes.size() << std::endl; + 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 +} diff --git a/sycl/test-e2e/DeviceImageBackendContent/basic_test.cpp b/sycl/test-e2e/DeviceImageBackendContent/basic_test.cpp new file mode 100644 index 0000000000000..8d0c873f1d2d3 --- /dev/null +++ b/sycl/test-e2e/DeviceImageBackendContent/basic_test.cpp @@ -0,0 +1,50 @@ +// RUN: %{build} -std=c++20 -o %t.out +// RUN: %{run} %t.out + +#include +#include +#include + +int main() { + sycl::queue q; + sycl::context ctxt = q.get_context(); + sycl::buffer buf(sycl::range<1>(1)); + 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); + 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 + std::span bytes_view; +#endif + for (const auto &img : bundle) { + // 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 + 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 + } + return 0; +} diff --git a/sycl/test/abi/sycl_symbols_linux.dump b/sycl/test/abi/sycl_symbols_linux.dump index 83edad971c869..144e77736de98 100644 --- a/sycl/test/abi/sycl_symbols_linux.dump +++ b/sycl/test/abi/sycl_symbols_linux.dump @@ -3742,6 +3742,8 @@ _ZNK4sycl3_V16detail16AccessorBaseHost6getPtrEv _ZNK4sycl3_V16detail16AccessorBaseHost9getOffsetEv _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_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 edc7bcd617733..e601a840b50d1 100644 --- a/sycl/test/abi/sycl_symbols_windows.dump +++ b/sycl/test/abi/sycl_symbols_windows.dump @@ -3862,6 +3862,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@@IEBA?AW4backend@34@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 diff --git a/sycl/test/extensions/DeviceImageBackendContent/negative_test.cpp b/sycl/test/extensions/DeviceImageBackendContent/negative_test.cpp new file mode 100644 index 0000000000000..bc5783e64ba64 --- /dev/null +++ b/sycl/test/extensions/DeviceImageBackendContent/negative_test.cpp @@ -0,0 +1,47 @@ +// RUN: %clang -fsycl -fsyntax-only -std=c++20 -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; +}