diff --git a/sycl/cmake/modules/FetchUnifiedRuntime.cmake b/sycl/cmake/modules/FetchUnifiedRuntime.cmake index a5a70600e9002..f8ec85942983c 100644 --- a/sycl/cmake/modules/FetchUnifiedRuntime.cmake +++ b/sycl/cmake/modules/FetchUnifiedRuntime.cmake @@ -116,14 +116,14 @@ if(SYCL_PI_UR_USE_FETCH_CONTENT) CACHE PATH "Path to external '${name}' adapter source dir" FORCE) endfunction() -set(UNIFIED_RUNTIME_REPO "https://github.com/oneapi-src/unified-runtime.git") - # commit 3e762e00bcf13d158fb58e8e8c2eabcfc8934b4e - # Merge: c805a71a a2a053de + set(UNIFIED_RUNTIME_REPO "https://github.com/oneapi-src/unified-runtime.git") + # commit c5d2175b5823d5b74de1e7e0d6081ab6d885bc34 + # Merge: 99489ad4 c86beb60 # Author: Omar Ahmed - # Date: Wed Jul 31 12:26:34 2024 +0100 - # Merge pull request #1884 from callumfare/callum/fix_printtrace - # Enable PrintTrace when SYCL UR tracing is enabled - set(UNIFIED_RUNTIME_TAG 3e762e00bcf13d158fb58e8e8c2eabcfc8934b4e) + # Date: Wed Jul 31 14:52:26 2024 +0100 + # Merge pull request #1882 from przemektmalon/przemek/interop-map-memory + # [Bindless][Exp] Add interop memory mapping to USM. + set(UNIFIED_RUNTIME_TAG c5d2175b5823d5b74de1e7e0d6081ab6d885bc34) set(UMF_BUILD_EXAMPLES OFF CACHE INTERNAL "EXAMPLES") # Due to the use of dependentloadflag and no installer for UMF and hwloc we need diff --git a/sycl/doc/extensions/experimental/sycl_ext_oneapi_bindless_images.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_oneapi_bindless_images.asciidoc index 1a7a9062885db..7218217298237 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_bindless_images.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_bindless_images.asciidoc @@ -1712,19 +1712,35 @@ image_mem_handle map_external_image_memory( external_mem externalMemHandle, const image_descriptor &imageDescriptor, const sycl::queue &syclQueue); + +void *map_external_linear_memory( + external_mem externalMemHandle, + uint64_t size, uint64_t offset, + const sycl::device &syclDevice, + const sycl::context &syclContext); +void *map_external_linear_memory( + external_mem externalMemHandle, + uint64_t size, uint64_t offset, + const sycl::queue &syclQueue); } ``` The resulting `external_mem` can then be mapped, where the resulting type -is an `image_mem_handle`. This can be used to construct images in the same way -as memory allocated through `alloc_image_mem`. The `ext_oneapi_copy` operations -also work with imported memory mapped to `image_mem_handle` types. - -When calling `create_image` with an `image_mem_handle` mapped from an external -memory object, the user must ensure that the image descriptor they pass to -`create_image` has members that match or map to those of the external API. -A mismatch between any of the `width`, `height`, `depth`, `image_channel_type`, -or `num_channels` members will result in undefined behavior. +is an `image_mem_handle` or a `void *`. This can be used to construct images in +the same way as memory allocated through `alloc_image_mem`, +`pitched_alloc_device`, or another USM allocation method. The `ext_oneapi_copy` +operations also work with imported memory mapped to `image_mem_handle` and +`void *` types. + +When calling `create_image` with an `image_mem_handle` or `void *` mapped from +an external memory object, the user must ensure that the image descriptor they +pass to `create_image` has members that match or map to those of the external +API. A mismatch between any of the `width`, `height`, `depth`, +`image_channel_type`, or `num_channels` members will result in undefined +behavior. Likewise, if the image is mapped to a linear USM (`void *`) region, +the pitch value passed to `create_image` needs to match the pitch of the image +as defined by the external API. Note that when external memory is mapped to a +linear USM region, this is device-side USM, and not accessible on the host. Additionally, the `image_type` describing the image must match to the image of the external API. The current supported importable image types are `standard` @@ -2866,4 +2882,6 @@ These features still need to be handled: - Removed `handle` keyword from `interop_xxx_handle` to clear up possible confusion between 3rd party interop handles and the imported `interop_xxx_handle`. +|5.17|2024-07-30| - Add support for mapping external memory to linear USM using + `map_external_linear_memory`. |====================== diff --git a/sycl/include/sycl/ext/oneapi/bindless_images.hpp b/sycl/include/sycl/ext/oneapi/bindless_images.hpp index dcd707aaa0be3..777ddecd887d5 100644 --- a/sycl/include/sycl/ext/oneapi/bindless_images.hpp +++ b/sycl/include/sycl/ext/oneapi/bindless_images.hpp @@ -171,7 +171,7 @@ image_mem_handle map_external_image_memory(external_mem extMem, const sycl::context &syclContext); /** - * @brief Maps an external memory handle to an image memory handle (which may + * @brief Maps an external memory object to an image memory handle (which may * have a device optimized memory layout) * * @param extMem External memory object @@ -184,6 +184,36 @@ image_mem_handle map_external_image_memory(external_mem extMem, const image_descriptor &desc, const sycl::queue &syclQueue); +/** + * @brief Maps an external memory object to a memory region described by the + * returned void * + * + * @param extMem External memory object + * @param offset Offset of memory region to map + * @param size Size of memory region to map + * @param syclDevice The device in which we create our image memory handle + * @param syclContext The context in which we create our image memory handle + * @return Memory handle to externally allocated memory on the device + */ +__SYCL_EXPORT +void *map_external_linear_memory(external_mem extMem, uint64_t offset, + uint64_t size, const sycl::device &syclDevice, + const sycl::context &syclContext); + +/** + * @brief Maps an external memory object to a memory region described by the + * returned void * + * + * @param extMem External memory object + * @param offset Offset of memory region to map + * @param size Size of memory region to map + * @param syclQueue The queue in which we create our image memory handle + * @return Memory handle to externally allocated memory on the device + */ +__SYCL_EXPORT +void *map_external_linear_memory(external_mem extMem, uint64_t offset, + uint64_t size, const sycl::queue &syclQueue); + /** * @brief Import external semaphore taking an external semaphore descriptor * (the type of which is dependent on the OS & external API) diff --git a/sycl/source/detail/bindless_images.cpp b/sycl/source/detail/bindless_images.cpp index 775dbd0bb60fd..457411313381f 100644 --- a/sycl/source/detail/bindless_images.cpp +++ b/sycl/source/detail/bindless_images.cpp @@ -556,6 +556,35 @@ image_mem_handle map_external_image_memory(external_mem extMem, syclQueue.get_context()); } +__SYCL_EXPORT +void *map_external_linear_memory(external_mem extMem, uint64_t offset, + uint64_t size, const sycl::device &syclDevice, + const sycl::context &syclContext) { + std::shared_ptr CtxImpl = + sycl::detail::getSyclObjImpl(syclContext); + ur_context_handle_t C = CtxImpl->getHandleRef(); + std::shared_ptr DevImpl = + sycl::detail::getSyclObjImpl(syclDevice); + ur_device_handle_t Device = DevImpl->getHandleRef(); + const sycl::detail::PluginPtr &Plugin = CtxImpl->getPlugin(); + + ur_exp_external_mem_handle_t urExternalMem{extMem.raw_handle}; + + void *retMemory; + Plugin->call(urBindlessImagesMapExternalLinearMemoryExp, + C, Device, offset, size, urExternalMem, + &retMemory); + + return retMemory; +} + +__SYCL_EXPORT +void *map_external_linear_memory(external_mem extMem, uint64_t offset, + uint64_t size, const sycl::queue &syclQueue) { + return map_external_linear_memory( + extMem, offset, size, syclQueue.get_device(), syclQueue.get_context()); +} + __SYCL_EXPORT void release_external_memory(external_mem extMem, const sycl::device &syclDevice, const sycl::context &syclContext) { diff --git a/sycl/test-e2e/bindless_images/vulkan_interop/sampled_images_USM.cpp b/sycl/test-e2e/bindless_images/vulkan_interop/sampled_images_USM.cpp new file mode 100644 index 0000000000000..a0b030cf3cb3a --- /dev/null +++ b/sycl/test-e2e/bindless_images/vulkan_interop/sampled_images_USM.cpp @@ -0,0 +1,355 @@ +// REQUIRES: cuda +// REQUIRES: vulkan + +// RUN: %{build} %link-vulkan -o %t.out +// RUN: %{run} %t.out + +#include "../helpers/common.hpp" +#include "vulkan_common.hpp" + +#include + +#include +#include + +namespace syclexp = sycl::ext::oneapi::experimental; + +struct handles_t { + syclexp::sampled_image_handle imgInput; + void *imgMem; + syclexp::external_mem inputInteropMemHandle; +}; + +template +handles_t create_test_handles(sycl::context &ctxt, sycl::device &dev, + const syclexp::bindless_image_sampler &samp, + InteropHandleT interopHandle, + syclexp::image_descriptor desc, + const size_t imgSize, + const size_t imgPitchBytes) { + // Extension: external memory descriptor +#ifdef _WIN32 + syclexp::external_mem_descriptor + inputExtMemDesc{interopHandle, + syclexp::external_mem_handle_type::win32_nt_handle, + imgSize}; +#else + syclexp::external_mem_descriptor inputExtMemDesc{ + interopHandle, syclexp::external_mem_handle_type::opaque_fd, imgSize}; +#endif + + // Extension: interop mem handle imported from file descriptor + syclexp::external_mem inputInteropMemHandle = + syclexp::import_external_memory(inputExtMemDesc, dev, ctxt); + + // Extension: interop mem handle imported from file descriptor + void *mappedImgMem = syclexp::map_external_linear_memory( + inputInteropMemHandle, 0, imgSize, dev, ctxt); + + // Extension: create the image and return the handle + syclexp::sampled_image_handle imgInput = + syclexp::create_image(mappedImgMem, imgPitchBytes, samp, desc, dev, ctxt); + + return {imgInput, mappedImgMem, inputInteropMemHandle}; +} + +template +bool run_sycl(InteropHandleT inputInteropMemHandle, size_t imgPitchBytes, + sycl::range globalSize, sycl::range localSize) { + sycl::device dev; + sycl::queue q(dev); + auto ctxt = q.get_context(); + + // Image descriptor - mapped to Vulkan image layout + syclexp::image_descriptor desc(globalSize, NChannels, CType); + + syclexp::bindless_image_sampler samp( + sycl::addressing_mode::repeat, + sycl::coordinate_normalization_mode::normalized, + sycl::filtering_mode::linear); + + const auto numElems = globalSize.size(); + + auto width = globalSize[0]; + auto height = globalSize[1]; + + const size_t img_size = + (imgPitchBytes / 8) * height * sizeof(DType) * NChannels; + + using VecType = sycl::vec; + + auto handles = create_test_handles(ctxt, dev, samp, inputInteropMemHandle, + desc, img_size, imgPitchBytes); + + sycl::range outBufferRange; + outBufferRange = sycl::range{height, width}; + + std::vector out(width * height); + try { + sycl::buffer buf((VecType *)out.data(), outBufferRange); + q.submit([&](sycl::handler &cgh) { + auto outAcc = buf.template get_access( + cgh, outBufferRange); + cgh.parallel_for( + sycl::nd_range{globalSize, localSize}, + [=](sycl::nd_item it) { + size_t dim0 = it.get_global_id(0); + size_t dim1 = it.get_global_id(1); + + // Normalize coordinates -- +0.5 to look towards centre of pixel + float fdim0 = float(dim0 + 0.5f) / (float)width; + float fdim1 = float(dim1 + 0.5f) / (float)height; + + // Extension: sample image data from handle (Vulkan imported) + VecType pixel = syclexp::sample_image( + handles.imgInput, sycl::float2(fdim0, fdim1)); + + pixel *= static_cast(10.1f); + outAcc[sycl::id{dim1, dim0}] = pixel; + }); + }); + q.wait_and_throw(); + + syclexp::destroy_image_handle(handles.imgInput, dev, ctxt); + sycl::free(handles.imgMem, ctxt); + syclexp::release_external_memory(handles.inputInteropMemHandle, dev, ctxt); + } catch (sycl::exception e) { + std::cerr << "\tKernel submission failed! " << e.what() << std::endl; + exit(-1); + } catch (...) { + std::cerr << "\tKernel submission failed!" << std::endl; + exit(-1); + } + + printString("Validating\n"); + bool validated = true; + + for (int i = 0; i < numElems; ++i) { + bool mismatch = false; + + VecType expected = bindless_helpers::init_vector(i) * + static_cast(10.1f); + if (!bindless_helpers::equal_vec(out[i], expected)) { + mismatch = true; + validated = false; + } + + if (mismatch) { +#ifdef VERBOSE_PRINT + std::cout << "Result mismatch! Expected: " << expected + << ", Actual: " << out[i] << "\n"; +#else + break; +#endif + } + } + + if (validated) { + printString("Results are correct!\n"); + } + + return validated; +} + +template +bool run_test(sycl::range dims, sycl::range localSize, + unsigned int seed = 0) { + uint32_t width = static_cast(dims[0]); + uint32_t height = 1; + uint32_t depth = 1; + + size_t numElems = dims[0]; + VkImageType imgType = VK_IMAGE_TYPE_1D; + + if constexpr (NDims > 1) { + numElems *= dims[1]; + height = static_cast(dims[1]); + imgType = VK_IMAGE_TYPE_2D; + } + + using VecType = sycl::vec; + + VkFormat format = vkutil::to_vulkan_format(COrder, CType); + + printString("Creating input image\n"); + // Create input image memory + auto inputImage = vkutil::createImage(imgType, format, {width, height, depth}, + VK_IMAGE_USAGE_TRANSFER_SRC_BIT | + VK_IMAGE_USAGE_TRANSFER_DST_BIT, + 1 /*mipLevels*/, true /*linearTiling*/); + + const size_t imageRowPitchBytes = + static_cast(vkutil::getImageRowPitch(inputImage)); + const size_t imageRowPitchElements = imageRowPitchBytes / 8; + const size_t imageSizeBytes = + imageRowPitchElements * height * depth * NChannels * sizeof(DType); + const size_t stagingBufferSizeBytes = imageSizeBytes; + + VkMemoryRequirements memRequirements = {}; + memRequirements.size = imageSizeBytes; + auto inputImageMemoryTypeIndex = vkutil::getImageMemoryTypeIndex( + inputImage, VK_MEMORY_PROPERTY_DEVICE_LOCAL_BIT, memRequirements); + auto inputMemory = + vkutil::allocateDeviceMemory(imageSizeBytes, inputImageMemoryTypeIndex); + VK_CHECK_CALL(vkBindImageMemory(vk_device, inputImage, inputMemory, + 0 /*memoryOffset*/)); + + printString("Creating staging buffers\n"); + // Create input staging memory + auto inputStagingBuffer = vkutil::createBuffer( + stagingBufferSizeBytes, + VK_BUFFER_USAGE_TRANSFER_SRC_BIT | VK_BUFFER_USAGE_TRANSFER_DST_BIT); + auto inputStagingMemoryTypeIndex = vkutil::getBufferMemoryTypeIndex( + inputStagingBuffer, VK_MEMORY_PROPERTY_HOST_VISIBLE_BIT | + VK_MEMORY_PROPERTY_HOST_COHERENT_BIT); + auto inputStagingMemory = vkutil::allocateDeviceMemory( + stagingBufferSizeBytes, inputStagingMemoryTypeIndex, + false /*exportable*/); + VK_CHECK_CALL(vkBindBufferMemory(vk_device, inputStagingBuffer, + inputStagingMemory, 0 /*memoryOffset*/)); + + printString("Populating staging buffer\n"); + // Populate staging memory + VecType *inputStagingData = nullptr; + VK_CHECK_CALL(vkMapMemory(vk_device, inputStagingMemory, 0 /*offset*/, + imageSizeBytes, 0 /*flags*/, + (void **)&inputStagingData)); + for (int j = 0; j < height; ++j) { + for (int i = 0; i < width; ++i) { + int index = (j * imageRowPitchElements) + i; + int element = (j * width) + i; + inputStagingData[index] = + bindless_helpers::init_vector(element); + } + } + vkUnmapMemory(vk_device, inputStagingMemory); + + printString("Submitting image layout transition\n"); + // Transition image layouts + { + VkImageMemoryBarrier barrierInput = + vkutil::createImageMemoryBarrier(inputImage, 1 /*mipLevels*/); + + VkCommandBufferBeginInfo cbbi = {}; + cbbi.sType = VK_STRUCTURE_TYPE_COMMAND_BUFFER_BEGIN_INFO; + cbbi.flags = VK_COMMAND_BUFFER_USAGE_ONE_TIME_SUBMIT_BIT; + + VK_CHECK_CALL(vkBeginCommandBuffer(vk_computeCmdBuffer, &cbbi)); + vkCmdPipelineBarrier(vk_computeCmdBuffer, VK_PIPELINE_STAGE_TOP_OF_PIPE_BIT, + VK_PIPELINE_STAGE_TRANSFER_BIT, 0, 0, nullptr, 0, + nullptr, 1, &barrierInput); + VK_CHECK_CALL(vkEndCommandBuffer(vk_computeCmdBuffer)); + + VkSubmitInfo submission = {}; + submission.sType = VK_STRUCTURE_TYPE_SUBMIT_INFO; + submission.commandBufferCount = 1; + submission.pCommandBuffers = &vk_computeCmdBuffer; + + VK_CHECK_CALL(vkQueueSubmit(vk_compute_queue, 1 /*submitCount*/, + &submission, VK_NULL_HANDLE /*fence*/)); + VK_CHECK_CALL(vkQueueWaitIdle(vk_compute_queue)); + } + + printString("Copying staging memory to images\n"); + // Copy staging to main image memory + { + VkCommandBufferBeginInfo cbbi = {}; + cbbi.sType = VK_STRUCTURE_TYPE_COMMAND_BUFFER_BEGIN_INFO; + cbbi.flags = VK_COMMAND_BUFFER_USAGE_ONE_TIME_SUBMIT_BIT; + + VkBufferImageCopy copyRegion = {}; + copyRegion.imageExtent = {width, height, depth}; + copyRegion.imageSubresource.aspectMask = VK_IMAGE_ASPECT_COLOR_BIT; + copyRegion.imageSubresource.layerCount = 1; + copyRegion.bufferRowLength = imageRowPitchElements; + + VK_CHECK_CALL(vkBeginCommandBuffer(vk_transferCmdBuffers[0], &cbbi)); + vkCmdCopyBufferToImage(vk_transferCmdBuffers[0], inputStagingBuffer, + inputImage, VK_IMAGE_LAYOUT_GENERAL, + 1 /*regionCount*/, ©Region); + VK_CHECK_CALL(vkEndCommandBuffer(vk_transferCmdBuffers[0])); + + VkSubmitInfo submission = {}; + submission.sType = VK_STRUCTURE_TYPE_SUBMIT_INFO; + submission.commandBufferCount = 1; + submission.pCommandBuffers = &vk_transferCmdBuffers[0]; + + VK_CHECK_CALL(vkQueueSubmit(vk_transfer_queue, 1 /*submitCount*/, + &submission, VK_NULL_HANDLE /*fence*/)); + VK_CHECK_CALL(vkQueueWaitIdle(vk_transfer_queue)); + } + + printString("Getting memory file descriptors\n"); + // Pass memory to SYCL for modification + +#ifdef _WIN32 + auto input_mem_handle = vkutil::getMemoryWin32Handle(inputMemory); +#else + auto input_mem_handle = vkutil::getMemoryOpaqueFD(inputMemory); +#endif + + printString("Calling into SYCL with interop memory handle\n"); + + bool validated = run_sycl( + input_mem_handle, imageRowPitchBytes, dims, localSize); + + // Cleanup + vkDestroyBuffer(vk_device, inputStagingBuffer, nullptr); + vkDestroyImage(vk_device, inputImage, nullptr); + vkFreeMemory(vk_device, inputStagingMemory, nullptr); + vkFreeMemory(vk_device, inputMemory, nullptr); + + return validated; +} + +bool run_tests() { + bool valid = run_test<2, float, 4, sycl::image_channel_type::fp32, + sycl::image_channel_order::rgba, class float_2d_4c>( + {16, 16}, {16, 16}, 0); + + valid &= run_test<2, float, 2, sycl::image_channel_type::fp32, + sycl::image_channel_order::rg, class float_2d_2c>( + {1024, 1024}, {32, 32}, 0); + + return valid; +} + +int main() { + + if (vkutil::setupInstance() != VK_SUCCESS) { + std::cerr << "Instance setup failed!\n"; + return EXIT_FAILURE; + } + + sycl::device dev; + + if (vkutil::setupDevice(dev.get_info()) != + VK_SUCCESS) { + std::cerr << "Device setup failed!\n"; + return EXIT_FAILURE; + } + + if (vkutil::setupCommandBuffers() != VK_SUCCESS) { + std::cerr << "Compute pipeline setup failed!\n"; + return EXIT_FAILURE; + } + + bool result_ok = run_tests(); + + if (vkutil::cleanup() != VK_SUCCESS) { + std::cerr << "Cleanup failed!\n"; + return EXIT_FAILURE; + } + + if (result_ok) { + std::cout << "All tests passed!\n"; + return EXIT_SUCCESS; + } + + std::cerr << "Test failed\n"; + return EXIT_FAILURE; +} diff --git a/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_common.hpp b/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_common.hpp index 096f11dd50369..4a70f2a8e9edc 100644 --- a/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_common.hpp +++ b/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_common.hpp @@ -472,7 +472,7 @@ program is compiled for. */ VkImage createImage(VkImageType type, VkFormat format, VkExtent3D extent, VkImageUsageFlags usage, size_t mipLevels, - bool exportable = true) { + bool linearTiling = false, bool exportable = true) { VkImageCreateInfo ici = {}; ici.sType = VK_STRUCTURE_TYPE_IMAGE_CREATE_INFO; ici.imageType = type; @@ -484,6 +484,10 @@ VkImage createImage(VkImageType type, VkFormat format, VkExtent3D extent, ici.sharingMode = VK_SHARING_MODE_EXCLUSIVE; ici.samples = VK_SAMPLE_COUNT_1_BIT; + if (linearTiling) { + ici.tiling = VK_IMAGE_TILING_LINEAR; + } + VkExternalMemoryImageCreateInfo emici = {}; if (exportable) { emici.sType = VK_STRUCTURE_TYPE_EXTERNAL_MEMORY_IMAGE_CREATE_INFO; @@ -504,6 +508,24 @@ VkImage createImage(VkImageType type, VkFormat format, VkExtent3D extent, return image; } +/* +Returns the row pitch with which a linear image's first subresource was +created with in bytes. +*/ +VkDeviceSize getImageRowPitch(VkImage image) { + + VkImageSubresource imageSubresource = {}; + imageSubresource.aspectMask = VK_IMAGE_ASPECT_COLOR_BIT; + imageSubresource.mipLevel = 0; + imageSubresource.arrayLayer = 0; + + VkSubresourceLayout subresourceLayout = {}; + vkGetImageSubresourceLayout(vk_device, image, &imageSubresource, + &subresourceLayout); + + return subresourceLayout.rowPitch; +} + /* Allocate `size` of device memory of the specified memory type. This function also allows users to specify whether the memory will be diff --git a/sycl/test/abi/sycl_symbols_linux.dump b/sycl/test/abi/sycl_symbols_linux.dump index 091fc56494f1d..0ab41c0c4019b 100644 --- a/sycl/test/abi/sycl_symbols_linux.dump +++ b/sycl/test/abi/sycl_symbols_linux.dump @@ -3045,6 +3045,8 @@ _ZN4sycl3_V13ext6oneapi12experimental25import_external_semaphoreINS3_21resource_ _ZN4sycl3_V13ext6oneapi12experimental25import_external_semaphoreINS3_21resource_win32_handleEEENS3_18external_semaphoreENS3_29external_semaphore_descriptorIT_EERKNS0_6deviceERKNS0_7contextE _ZN4sycl3_V13ext6oneapi12experimental25map_external_image_memoryENS3_12external_memERKNS3_16image_descriptorERKNS0_5queueE _ZN4sycl3_V13ext6oneapi12experimental25map_external_image_memoryENS3_12external_memERKNS3_16image_descriptorERKNS0_6deviceERKNS0_7contextE +_ZN4sycl3_V13ext6oneapi12experimental26map_external_linear_memoryENS3_12external_memEmmRKNS0_5queueE +_ZN4sycl3_V13ext6oneapi12experimental26map_external_linear_memoryENS3_12external_memEmmRKNS0_6deviceERKNS0_7contextE _ZN4sycl3_V13ext6oneapi12experimental26release_external_semaphoreENS3_18external_semaphoreERKNS0_5queueE _ZN4sycl3_V13ext6oneapi12experimental26release_external_semaphoreENS3_18external_semaphoreERKNS0_6deviceERKNS0_7contextE _ZN4sycl3_V13ext6oneapi12experimental4node12update_rangeILi1EEEvNS0_5rangeIXT_EEE diff --git a/sycl/test/abi/sycl_symbols_windows.dump b/sycl/test/abi/sycl_symbols_windows.dump index 72fe0ebad1ed8..cc50e445d829b 100644 --- a/sycl/test/abi/sycl_symbols_windows.dump +++ b/sycl/test/abi/sycl_symbols_windows.dump @@ -4128,6 +4128,8 @@ ?map@physical_mem@experimental@oneapi@ext@_V1@sycl@@QEBAPEAX_K0W4address_access_mode@23456@0@Z ?map_external_image_memory@experimental@oneapi@ext@_V1@sycl@@YA?AUimage_mem_handle@12345@Uexternal_mem@12345@AEBUimage_descriptor@12345@AEBVdevice@45@AEBVcontext@45@@Z ?map_external_image_memory@experimental@oneapi@ext@_V1@sycl@@YA?AUimage_mem_handle@12345@Uexternal_mem@12345@AEBUimage_descriptor@12345@AEBVqueue@45@@Z +?map_external_linear_memory@experimental@oneapi@ext@_V1@sycl@@YAPEAXUexternal_mem@12345@_K1AEBVdevice@45@AEBVcontext@45@@Z +?map_external_linear_memory@experimental@oneapi@ext@_V1@sycl@@YAPEAXUexternal_mem@12345@_K1AEBVqueue@45@@Z ?markBufferAsInternal@detail@_V1@sycl@@YAXAEBV?$shared_ptr@Vbuffer_impl@detail@_V1@sycl@@@std@@@Z ?mem_advise@experimental@oneapi@ext@_V1@sycl@@YAXVqueue@45@PEAX_KHAEBUcode_location@detail@45@@Z ?mem_advise@handler@_V1@sycl@@QEAAXPEBX_KH@Z