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 5c63e54a5b744..92b17f5a7f0bb 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_bindless_images.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_bindless_images.asciidoc @@ -2150,6 +2150,10 @@ the same way as memory allocated through `alloc_image_mem`, operations also work with imported memory mapped to `image_mem_handle` and `void *` types. +An `external_mem` handle can only be mapped to a single `image_mem_handle` or +`void *` at any one time. Attempting to map a single `external_mem` handle to +more than one `image_mem_handle` or `void *` is considered undefined behaviour. + 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 @@ -2165,11 +2169,41 @@ the external API. The current supported importable image types are `standard` and `mipmap`. Attempting to import other image types will result in undefined behaviour. -Once a user has finished operating on imported memory, they must ensure that -they destroy the imported memory handle through `release_external_memory`. +Once a user has finished operating on mapped memory, they must ensure that they +unmap that memory. + +Memory mapped using `map_external_image_memory` should be unmapped using +`unmap_external_image_memory`. The `image_type` parameter passed to this +function must reflect the `image_type` of the image descriptor used when the +memory was originally mapped. Passing an `image_type` value different to that of +the value used in the image descriptor when the memory was originally mapped +will result in undefined behaviour. -`release_external_memory` can only accept `external_mem` objects that were -created through `import_external_memory`. +Memory mapped using `map_external_linear_memory` should be unmapped using +`unmap_external_linear_memory`. + +```cpp +namespace sycl::ext::oneapi::experimental { + +void unmap_external_image_memory(image_mem_handle mappedImageMem, + image_type imageType, + const sycl::device &syclDevice, + const sycl::context &syclContext); +void unmap_external_image_memory(image_mem_handle mappedImageMem, + image_type imageType, + const sycl::queue &syclQueue); + +void unmap_external_linear_memory(void *mappedLinearMem, + const sycl::device &syclDevice, + const sycl::context &syclContext); +void unmap_external_linear_memory(void *mappedLinearMem, + const sycl::queue &syclQueue); +} +``` + +Once all memory mapped from a given `external_mem` handle has been unmapped, +and the user has finished operating on the external memory, they should then +release the `externa_mem` handle using `release_external_memory`. ```cpp namespace sycl::ext::oneapi::experimental { @@ -2182,9 +2216,6 @@ void release_external_memory(external_mem externalMem, } ``` -Destroying or freeing any imported memory through `image_mem_free` or -`sycl::free` will result in undefined behavior. - === Importing external semaphores [[importing_external_semaphores]] In addition to proposing importation of external memory resources, we also diff --git a/sycl/include/sycl/ext/oneapi/bindless_images.hpp b/sycl/include/sycl/ext/oneapi/bindless_images.hpp index f49a4c1473260..9c5da0e8650f4 100644 --- a/sycl/include/sycl/ext/oneapi/bindless_images.hpp +++ b/sycl/include/sycl/ext/oneapi/bindless_images.hpp @@ -285,6 +285,54 @@ __SYCL_EXPORT void release_external_memory(external_mem externalMem, __SYCL_EXPORT void release_external_memory(external_mem externalMem, const sycl::queue &syclQueue); +/** + * @brief Unmap external linear memory region + * + * @param mappedLinearMem Pointer to the mapped linear memory region to unmap + * @param syclDevice The device in which the external memory was created + * @param syclContext The context in which the external memory was created + */ +__SYCL_EXPORT void +unmap_external_linear_memory(void *mappedLinearMem, + const sycl::device &syclDevice, + const sycl::context &syclContext); + +/** + * @brief Unmap external linear memory region + * + * @param mappedLinearMem Pointer to the mapped linear memory region to unmap + * @param syclQueue The queue in which the external memory was created + */ +inline void unmap_external_linear_memory(void *mappedLinearMem, + const sycl::queue &syclQueue) { + unmap_external_linear_memory(mappedLinearMem, syclQueue.get_device(), + syclQueue.get_context()); +} + +/** + * @brief Unmap external image memory + * + * @param mappedImageMem Handle to the mapped image memory to unmap + * @param syclDevice The device in which the external memory was created + * @param syclContext The context in which the external memory was created + */ +__SYCL_EXPORT void unmap_external_image_memory( + image_mem_handle mappedImageMem, image_type imageType, + const sycl::device &syclDevice, const sycl::context &syclContext); + +/** + * @brief Unmap external image memory + * + * @param mappedImageMem Handle to the mapped image memory to unmap + * @param syclQueue The queue in which the external memory was created + */ +inline void unmap_external_image_memory(image_mem_handle mappedImageMem, + image_type imageType, + const sycl::queue &syclQueue) { + unmap_external_image_memory(mappedImageMem, imageType, syclQueue.get_device(), + syclQueue.get_context()); +} + /** * @brief Create an image and return the device image handle * diff --git a/sycl/source/detail/bindless_images.cpp b/sycl/source/detail/bindless_images.cpp index fd4d812854851..e23df95129fbd 100644 --- a/sycl/source/detail/bindless_images.cpp +++ b/sycl/source/detail/bindless_images.cpp @@ -546,6 +546,24 @@ __SYCL_EXPORT void release_external_memory(external_mem extMem, syclQueue.get_context()); } +__SYCL_EXPORT void +unmap_external_linear_memory(void *mappedLinearRegion, + const sycl::device &syclDevice, + const sycl::context &syclContext) { + auto [urDevice, urCtx, Adapter] = get_ur_handles(syclDevice, syclContext); + + Adapter->call< + sycl::errc::invalid, + sycl::detail::UrApiKind::urBindlessImagesFreeMappedLinearMemoryExp>( + urCtx, urDevice, mappedLinearRegion); +} + +__SYCL_EXPORT void unmap_external_image_memory( + image_mem_handle mappedImageMem, image_type imageType, + const sycl::device &syclDevice, const sycl::context &syclContext) { + free_image_mem(mappedImageMem, imageType, syclDevice, syclContext); +} + template <> __SYCL_EXPORT external_semaphore import_external_semaphore( external_semaphore_descriptor externalSemaphoreDesc, diff --git a/sycl/test-e2e/bindless_images/vulkan_interop/buffer_usm.cpp b/sycl/test-e2e/bindless_images/vulkan_interop/buffer_usm.cpp new file mode 100644 index 0000000000000..0e9e956bbd683 --- /dev/null +++ b/sycl/test-e2e/bindless_images/vulkan_interop/buffer_usm.cpp @@ -0,0 +1,327 @@ +// REQUIRES: aspect-ext_oneapi_external_memory_import +// REQUIRES: vulkan + +// RUN: %{build} %link-vulkan -o %t.out %if target-spir %{ -Wno-ignored-attributes %} +// RUN: %{run} env NEOReadDebugKeys=1 UseBindlessMode=1 UseExternalAllocatorForSshAndDsh=1 %t.out + +/** + * This test does not use any image specific APIs. + * + * It is only testing the exportation of VkBuffer memory, and importing and + * mapping the VkBuffer memory into SYCL device memory. The imported SYCL device + * memory is then manipulated purely through pointers inside the kernel. + */ + +// Uncomment to print additional test information +// #define VERBOSE_PRINT + +#include "../helpers/common.hpp" +#include "vulkan_common.hpp" + +#include +#include + +namespace syclexp = sycl::ext::oneapi::experimental; + +template +void runSycl(const sycl::device &syclDevice, sycl::range<1> globalSize, + sycl::range<1> localSize, InteropMemHandleT extMemInHandle, + InteropMemHandleT extMemOutHandle) { + + sycl::queue syclQueue{syclDevice}; + + const size_t bufferSizeBytes = globalSize.size() * sizeof(uint32_t); + +#ifdef _WIN32 + syclexp::external_mem_descriptor extMemInDesc{ + extMemInHandle, syclexp::external_mem_handle_type::win32_nt_handle, + bufferSizeBytes}; + syclexp::external_mem_descriptor + extMemOutDesc{extMemOutHandle, + syclexp::external_mem_handle_type::win32_nt_handle, + bufferSizeBytes}; +#else + syclexp::external_mem_descriptor extMemInDesc{ + extMemInHandle, syclexp::external_mem_handle_type::opaque_fd, + bufferSizeBytes}; + syclexp::external_mem_descriptor extMemOutDesc{ + extMemOutHandle, syclexp::external_mem_handle_type::opaque_fd, + bufferSizeBytes}; +#endif + + // Extension: create interop memory handles. + syclexp::external_mem externalMemIn = + syclexp::import_external_memory(extMemInDesc, syclQueue); + syclexp::external_mem externalMemOut = + syclexp::import_external_memory(extMemOutDesc, syclQueue); + + // Extension: map linear memory handles. + uint32_t *memIn = static_cast(syclexp::map_external_linear_memory( + externalMemIn, 0 /* offset */, bufferSizeBytes, syclQueue)); + uint32_t *memOut = + static_cast(syclexp::map_external_linear_memory( + externalMemOut, 0 /* offset */, bufferSizeBytes, syclQueue)); + + try { + syclQueue.submit([&](sycl::handler &cgh) { + cgh.parallel_for( + sycl::nd_range<1>{globalSize, localSize}, [=](sycl::nd_item<1> it) { + size_t index = it.get_global_id(0); + + uint32_t bufferValue = memIn[index]; + memOut[index] = bufferValue * 2; + }); + }); + + // Wait for kernel completion before destroying external objects. + syclQueue.wait_and_throw(); + + // Cleanup. + syclexp::unmap_external_linear_memory(memIn, syclQueue); + syclexp::unmap_external_linear_memory(memOut, syclQueue); + syclexp::release_external_memory(externalMemIn, syclQueue); + syclexp::release_external_memory(externalMemOut, syclQueue); + + } 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); + } +} + +bool runTest(const sycl::device &syclDevice, sycl::range<1> bufferSize, + sycl::range<1> localSize) { + const size_t bufferSizeElems = bufferSize[0]; + const size_t bufferSizeBytes = bufferSizeElems * sizeof(uint32_t); + + VkBuffer vkInputBuffer; + VkDeviceMemory vkInputBufferMemory; + VkBuffer vkOutputBuffer; + VkDeviceMemory vkOutputBufferMemory; + + // Initialize buffer input data. + std::vector inputVec(bufferSizeElems, 0.f); + for (uint32_t i = 0; i < bufferSizeElems; ++i) { + inputVec[i] = i; + } + + // Create/allocate device buffers. + { + vkInputBuffer = vkutil::createBuffer(bufferSizeBytes, + VK_BUFFER_USAGE_TRANSFER_SRC_BIT | + VK_BUFFER_USAGE_TRANSFER_DST_BIT | + VK_BUFFER_USAGE_STORAGE_BUFFER_BIT, + true /*exportable*/); + auto inputBufferMemTypeIndex = vkutil::getBufferMemoryTypeIndex( + vkInputBuffer, VK_MEMORY_PROPERTY_DEVICE_LOCAL_BIT); + vkInputBufferMemory = vkutil::allocateDeviceMemory( + bufferSizeBytes, inputBufferMemTypeIndex, VK_NULL_HANDLE /*image*/); + VK_CHECK_CALL(vkBindBufferMemory(vk_device, vkInputBuffer, + vkInputBufferMemory, 0 /*memoryOffset*/)); + + vkOutputBuffer = vkutil::createBuffer( + bufferSizeBytes, + VK_BUFFER_USAGE_TRANSFER_SRC_BIT | VK_BUFFER_USAGE_TRANSFER_DST_BIT | + VK_BUFFER_USAGE_STORAGE_BUFFER_BIT, + true /* exportable */); + auto outputBufferMemTypeIndex = vkutil::getBufferMemoryTypeIndex( + vkOutputBuffer, VK_MEMORY_PROPERTY_DEVICE_LOCAL_BIT); + vkOutputBufferMemory = vkutil::allocateDeviceMemory( + bufferSizeBytes, outputBufferMemTypeIndex, VK_NULL_HANDLE /*image*/); + VK_CHECK_CALL(vkBindBufferMemory(vk_device, vkOutputBuffer, + vkOutputBufferMemory, 0 /*memoryOffset*/)); + } + + // Allocate temporary staging buffer and copy input data to device. + printString("Allocating staging memory and copying to device buffer\n"); + { + VkBuffer stagingBuffer; + VkDeviceMemory stagingMemory; + + stagingBuffer = vkutil::createBuffer(bufferSizeBytes, + VK_BUFFER_USAGE_TRANSFER_SRC_BIT | + VK_BUFFER_USAGE_TRANSFER_DST_BIT); + auto inputStagingMemTypeIndex = vkutil::getBufferMemoryTypeIndex( + stagingBuffer, VK_MEMORY_PROPERTY_HOST_VISIBLE_BIT | + VK_MEMORY_PROPERTY_HOST_COHERENT_BIT); + stagingMemory = vkutil::allocateDeviceMemory( + bufferSizeBytes, inputStagingMemTypeIndex, VK_NULL_HANDLE /*image*/, + false /*exportable*/); + VK_CHECK_CALL(vkBindBufferMemory(vk_device, stagingBuffer, stagingMemory, + 0 /*memoryOffset*/)); + + // Copy host data to temporary staging buffer. + uint32_t *inputStagingData = nullptr; + VK_CHECK_CALL(vkMapMemory(vk_device, stagingMemory, 0 /*offset*/, + bufferSizeBytes, 0 /*flags*/, + (void **)&inputStagingData)); + for (int i = 0; i < bufferSizeElems; ++i) { + inputStagingData[i] = inputVec[i]; + } + vkUnmapMemory(vk_device, stagingMemory); + + // Copy temporary staging buffer to device local buffer. + VkCommandBufferBeginInfo cbbi = {}; + cbbi.sType = VK_STRUCTURE_TYPE_COMMAND_BUFFER_BEGIN_INFO; + cbbi.flags = VK_COMMAND_BUFFER_USAGE_ONE_TIME_SUBMIT_BIT; + + VkBufferCopy copyRegion = {}; + copyRegion.size = bufferSizeBytes; + + VK_CHECK_CALL(vkBeginCommandBuffer(vk_transferCmdBuffers[0], &cbbi)); + vkCmdCopyBuffer(vk_transferCmdBuffers[0], stagingBuffer, vkInputBuffer, + 1 /*regionCount*/, ©Region); + VK_CHECK_CALL(vkEndCommandBuffer(vk_transferCmdBuffers[0])); + + std::vector stages{VK_PIPELINE_STAGE_TOP_OF_PIPE_BIT}; + + VkSubmitInfo submission = {}; + submission.sType = VK_STRUCTURE_TYPE_SUBMIT_INFO; + submission.commandBufferCount = 1; + submission.pCommandBuffers = &vk_transferCmdBuffers[0]; + submission.pWaitDstStageMask = stages.data(); + + VK_CHECK_CALL(vkQueueSubmit(vk_transfer_queue, 1 /*submitCount*/, + &submission, VK_NULL_HANDLE /*fence*/)); + VK_CHECK_CALL(vkQueueWaitIdle(vk_transfer_queue)); + + // Destroy temporary staging buffer and free memory. + vkDestroyBuffer(vk_device, stagingBuffer, nullptr); + vkFreeMemory(vk_device, stagingMemory, nullptr); + } + + printString("Getting memory interop handles\n"); + // Get memory interop handles. +#ifdef _WIN32 + auto bufferMemIn = vkutil::getMemoryWin32Handle(vkInputBufferMemory); + auto bufferMemOut = vkutil::getMemoryWin32Handle(vkOutputBufferMemory); +#else + auto bufferMemIn = vkutil::getMemoryOpaqueFD(vkInputBufferMemory); + auto bufferMemOut = vkutil::getMemoryOpaqueFD(vkOutputBufferMemory); +#endif + + // Call into SYCL to read from input buffer, and populate the output buffer. + printString("Calling into SYCL with interop memory handles\n"); + runSycl(syclDevice, bufferSize, localSize, bufferMemIn, bufferMemOut); + + // Copy device buffer memory to temporary staging buffer, and back to host. + printString("Copying buffer memory to host\n"); + std::vector outputVec(bufferSizeElems, 0); + { + VkBuffer stagingBuffer; + VkDeviceMemory stagingMemory; + + stagingBuffer = vkutil::createBuffer(bufferSizeBytes, + VK_BUFFER_USAGE_TRANSFER_SRC_BIT | + VK_BUFFER_USAGE_TRANSFER_DST_BIT); + auto outputStagingMemoryTypeIndex = vkutil::getBufferMemoryTypeIndex( + stagingBuffer, VK_MEMORY_PROPERTY_HOST_VISIBLE_BIT | + VK_MEMORY_PROPERTY_HOST_COHERENT_BIT); + stagingMemory = vkutil::allocateDeviceMemory( + bufferSizeBytes, outputStagingMemoryTypeIndex, VK_NULL_HANDLE /*image*/, + false /*exportable*/); + VK_CHECK_CALL(vkBindBufferMemory(vk_device, stagingBuffer, stagingMemory, + 0 /*memoryOffset*/)); + + VkCommandBufferBeginInfo cbbi = {}; + cbbi.sType = VK_STRUCTURE_TYPE_COMMAND_BUFFER_BEGIN_INFO; + cbbi.flags = VK_COMMAND_BUFFER_USAGE_ONE_TIME_SUBMIT_BIT; + + VkBufferCopy copyRegion = {}; + copyRegion.size = bufferSizeBytes; + + VK_CHECK_CALL(vkBeginCommandBuffer(vk_transferCmdBuffers[1], &cbbi)); + vkCmdCopyBuffer(vk_transferCmdBuffers[1], vkOutputBuffer, stagingBuffer, + 1 /*regionCount*/, ©Region); + VK_CHECK_CALL(vkEndCommandBuffer(vk_transferCmdBuffers[1])); + + std::vector stages{VK_PIPELINE_STAGE_TOP_OF_PIPE_BIT}; + + VkSubmitInfo submission = {}; + submission.sType = VK_STRUCTURE_TYPE_SUBMIT_INFO; + submission.commandBufferCount = 1; + submission.pCommandBuffers = &vk_transferCmdBuffers[1]; + submission.pWaitDstStageMask = stages.data(); + + VK_CHECK_CALL(vkQueueSubmit(vk_transfer_queue, 1 /*submitCount*/, + &submission, VK_NULL_HANDLE /*fence*/)); + VK_CHECK_CALL(vkQueueWaitIdle(vk_transfer_queue)); + + // Copy temporary staging buffer output data to host output vector. + uint32_t *outputStagingData = (uint32_t *)outputVec.data(); + VK_CHECK_CALL(vkMapMemory(vk_device, stagingMemory, 0 /*offset*/, + bufferSizeBytes, 0 /*flags*/, + (void **)&outputStagingData)); + for (int i = 0; i < bufferSizeElems; ++i) { + outputVec[i] = outputStagingData[i]; + } + vkUnmapMemory(vk_device, stagingMemory); + + // Destroy temporary staging buffer and free memory. + vkDestroyBuffer(vk_device, stagingBuffer, nullptr); + vkFreeMemory(vk_device, stagingMemory, nullptr); + } + + // Destroy buffers and free their memory. + vkDestroyBuffer(vk_device, vkInputBuffer, nullptr); + vkDestroyBuffer(vk_device, vkOutputBuffer, nullptr); + vkFreeMemory(vk_device, vkInputBufferMemory, nullptr); + vkFreeMemory(vk_device, vkOutputBufferMemory, nullptr); + + // Validate that SYCL made changes to the memory. + bool validated = true; + for (int i = 0; i < bufferSizeElems; ++i) { + uint32_t expected = inputVec[i] * 2; + if (outputVec[i] != expected) { + std::cerr << "Result mismatch! actual[" << i << "] == " << outputVec[i] + << " : expected == " << expected << "\n"; + validated = false; + } + if (!validated) + break; + } + + if (validated) { + printString("Results are correct!\n"); + } + + return validated; +} + +int main() { + + if (vkutil::setupInstance() != VK_SUCCESS) { + std::cerr << "Instance setup failed!\n"; + return EXIT_FAILURE; + } + + sycl::device syclDevice; + + if (vkutil::setupDevice(syclDevice.get_info()) != + VK_SUCCESS) { + std::cerr << "Device setup failed!\n"; + return EXIT_FAILURE; + } + + if (vkutil::setupCommandBuffers() != VK_SUCCESS) { + std::cerr << "Command buffers setup failed!\n"; + return EXIT_FAILURE; + } + + auto testPassed = runTest(syclDevice, {1024}, {256}); + + if (vkutil::cleanup() != VK_SUCCESS) { + std::cerr << "Cleanup failed!\n"; + return EXIT_FAILURE; + } + + if (testPassed) { + std::cout << "Test passed!\n"; + return EXIT_SUCCESS; + } + + std::cerr << "Test failed\n"; + return EXIT_FAILURE; +} diff --git a/sycl/test-e2e/bindless_images/vulkan_interop/depth_format.cpp b/sycl/test-e2e/bindless_images/vulkan_interop/depth_format.cpp index ccfe1119927ce..bb83addd5b0af 100644 --- a/sycl/test-e2e/bindless_images/vulkan_interop/depth_format.cpp +++ b/sycl/test-e2e/bindless_images/vulkan_interop/depth_format.cpp @@ -83,7 +83,8 @@ void runSycl(const sycl::device &syclDevice, sycl::range<2> globalSize, // Cleanup. syclexp::destroy_image_handle(imgIn, syclQueue); syclexp::destroy_image_handle(imgOut, syclQueue); - syclexp::free_image_mem(imgMemIn, syclexp::image_type::standard, syclQueue); + syclexp::unmap_external_image_memory( + imgMemIn, syclexp::image_type::standard, syclQueue); syclexp::free_image_mem(imgMemOut, syclexp::image_type::standard, syclQueue); syclexp::release_external_memory(externalMemIn, syclQueue); diff --git a/sycl/test-e2e/bindless_images/vulkan_interop/mipmaps.cpp b/sycl/test-e2e/bindless_images/vulkan_interop/mipmaps.cpp index 1a9d1275cc947..72d4da3a597b8 100644 --- a/sycl/test-e2e/bindless_images/vulkan_interop/mipmaps.cpp +++ b/sycl/test-e2e/bindless_images/vulkan_interop/mipmaps.cpp @@ -150,8 +150,8 @@ bool run_sycl(sycl::range globalSize, sycl::range localSize, q.wait_and_throw(); syclexp::destroy_image_handle(handles.imgInput, dev, ctxt); - syclexp::free_image_mem(handles.imgMem, syclexp::image_type::mipmap, dev, - ctxt); + syclexp::unmap_external_image_memory( + handles.imgMem, syclexp::image_type::mipmap, dev, ctxt); syclexp::release_external_memory(handles.inputExternalMem, dev, ctxt); } catch (sycl::exception e) { std::cerr << "\tKernel submission failed! " << e.what() << std::endl; diff --git a/sycl/test-e2e/bindless_images/vulkan_interop/sampled_images.cpp b/sycl/test-e2e/bindless_images/vulkan_interop/sampled_images.cpp index 46b34eb6d482a..95988df3a00fd 100644 --- a/sycl/test-e2e/bindless_images/vulkan_interop/sampled_images.cpp +++ b/sycl/test-e2e/bindless_images/vulkan_interop/sampled_images.cpp @@ -203,8 +203,8 @@ bool run_sycl(sycl::queue syclQueue, sycl::range globalSize, dev, ctxt); #endif syclexp::destroy_image_handle(handles.imgInput, dev, ctxt); - syclexp::free_image_mem(handles.imgMem, syclexp::image_type::standard, dev, - ctxt); + syclexp::unmap_external_image_memory( + handles.imgMem, syclexp::image_type::standard, dev, ctxt); syclexp::release_external_memory(handles.inputExternalMem, dev, ctxt); } catch (sycl::exception e) { std::cerr << "\tKernel submission failed! " << e.what() << std::endl; 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 index 702cddd626e70..562f27e9189e9 100644 --- a/sycl/test-e2e/bindless_images/vulkan_interop/sampled_images_USM.cpp +++ b/sycl/test-e2e/bindless_images/vulkan_interop/sampled_images_USM.cpp @@ -113,7 +113,7 @@ bool run_sycl(InteropHandleT inputInteropMemHandle, size_t imgPitchBytes, q.wait_and_throw(); syclexp::destroy_image_handle(handles.imgInput, dev, ctxt); - sycl::free(handles.imgMem, ctxt); + syclexp::unmap_external_linear_memory(handles.imgMem, dev, ctxt); syclexp::release_external_memory(handles.inputInteropMemHandle, dev, ctxt); } catch (sycl::exception e) { std::cerr << "\tKernel submission failed! " << e.what() << std::endl; diff --git a/sycl/test-e2e/bindless_images/vulkan_interop/unsampled_images.cpp b/sycl/test-e2e/bindless_images/vulkan_interop/unsampled_images.cpp index c4e200f466313..aee37c6477b91 100644 --- a/sycl/test-e2e/bindless_images/vulkan_interop/unsampled_images.cpp +++ b/sycl/test-e2e/bindless_images/vulkan_interop/unsampled_images.cpp @@ -145,12 +145,12 @@ void cleanup_test(sycl::context &ctxt, sycl::device &dev, handles_t handles) { syclexp::destroy_image_handle(handles.input_1, dev, ctxt); syclexp::destroy_image_handle(handles.input_2, dev, ctxt); syclexp::destroy_image_handle(handles.output, dev, ctxt); - syclexp::free_image_mem(handles.input_mem_handle_1, - syclexp::image_type::standard, dev, ctxt); - syclexp::free_image_mem(handles.input_mem_handle_2, - syclexp::image_type::standard, dev, ctxt); - syclexp::free_image_mem(handles.output_mem_handle, - syclexp::image_type::standard, dev, ctxt); + syclexp::unmap_external_image_memory( + handles.input_mem_handle_1, syclexp::image_type::standard, dev, ctxt); + syclexp::unmap_external_image_memory( + handles.input_mem_handle_2, syclexp::image_type::standard, dev, ctxt); + syclexp::unmap_external_image_memory( + handles.output_mem_handle, syclexp::image_type::standard, dev, ctxt); syclexp::release_external_memory(handles.input_external_mem_1, dev, ctxt); syclexp::release_external_memory(handles.input_external_mem_2, dev, ctxt); syclexp::release_external_memory(handles.output_external_mem, dev, ctxt); diff --git a/sycl/test-e2e/bindless_images/vulkan_interop/unsampled_images_timeline_semaphore.cpp b/sycl/test-e2e/bindless_images/vulkan_interop/unsampled_images_timeline_semaphore.cpp index e387bb7db126f..24cc79ef59840 100644 --- a/sycl/test-e2e/bindless_images/vulkan_interop/unsampled_images_timeline_semaphore.cpp +++ b/sycl/test-e2e/bindless_images/vulkan_interop/unsampled_images_timeline_semaphore.cpp @@ -122,10 +122,10 @@ void cleanup_sycl(sycl::context &ctxt, sycl::device &dev, handles_t handles) { ctxt); syclexp::destroy_image_handle(handles.input, dev, ctxt); syclexp::destroy_image_handle(handles.output, dev, ctxt); - syclexp::free_image_mem(handles.inputMemHandle, syclexp::image_type::standard, - dev, ctxt); - syclexp::free_image_mem(handles.outputMemHandle, - syclexp::image_type::standard, dev, ctxt); + syclexp::unmap_external_image_memory( + handles.inputMemHandle, syclexp::image_type::standard, dev, ctxt); + syclexp::unmap_external_image_memory( + handles.outputMemHandle, syclexp::image_type::standard, dev, ctxt); syclexp::release_external_memory(handles.inputExternalMem, dev, ctxt); syclexp::release_external_memory(handles.outputExternalMem, dev, ctxt); } 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 ee184e139f969..80ffb6fc9d832 100644 --- a/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_common.hpp +++ b/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_common.hpp @@ -461,13 +461,25 @@ VkResult setupCommandBuffers() { /* Create a Vulkan buffer with a specified size and usage. */ -VkBuffer createBuffer(size_t size, VkBufferUsageFlags usage) { +VkBuffer createBuffer(size_t size, VkBufferUsageFlags usage, + bool exportable = false) { VkBufferCreateInfo bci = {}; bci.sType = VK_STRUCTURE_TYPE_BUFFER_CREATE_INFO; bci.size = size; bci.usage = usage; bci.sharingMode = VK_SHARING_MODE_EXCLUSIVE; + VkExternalMemoryBufferCreateInfo embci = {}; + if (exportable) { + embci.sType = VK_STRUCTURE_TYPE_EXTERNAL_MEMORY_BUFFER_CREATE_INFO; +#ifdef _WIN32 + embci.handleTypes = VK_EXTERNAL_MEMORY_HANDLE_TYPE_OPAQUE_WIN32_BIT; +#else + embci.handleTypes = VK_EXTERNAL_MEMORY_HANDLE_TYPE_OPAQUE_FD_BIT; +#endif + bci.pNext = &embci; + } + VkBuffer buffer; if (vkCreateBuffer(vk_device, &bci, nullptr, &buffer) != VK_SUCCESS) { std::cerr << "Could not create buffer!\n"; diff --git a/sycl/test/abi/sycl_symbols_linux.dump b/sycl/test/abi/sycl_symbols_linux.dump index 232cc56b7f039..1e1ca428eff80 100644 --- a/sycl/test/abi/sycl_symbols_linux.dump +++ b/sycl/test/abi/sycl_symbols_linux.dump @@ -3067,6 +3067,8 @@ _ZN4sycl3_V13ext6oneapi12experimental26map_external_linear_memoryENS3_12external _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_V13ext6oneapi12experimental27unmap_external_image_memoryENS3_16image_mem_handleENS3_10image_typeERKNS0_6deviceERKNS0_7contextE +_ZN4sycl3_V13ext6oneapi12experimental28unmap_external_linear_memoryEPvRKNS0_6deviceERKNS0_7contextE _ZN4sycl3_V13ext6oneapi12experimental4node12update_rangeILi1EEEvNS0_5rangeIXT_EEE _ZN4sycl3_V13ext6oneapi12experimental4node12update_rangeILi2EEEvNS0_5rangeIXT_EEE _ZN4sycl3_V13ext6oneapi12experimental4node12update_rangeILi3EEEvNS0_5rangeIXT_EEE diff --git a/sycl/test/abi/sycl_symbols_windows.dump b/sycl/test/abi/sycl_symbols_windows.dump index cbd3ae8299c59..13c10fe74d305 100644 --- a/sycl/test/abi/sycl_symbols_windows.dump +++ b/sycl/test/abi/sycl_symbols_windows.dump @@ -4437,6 +4437,8 @@ ?throwOnKernelParameterMisuseHelper@handler@_V1@sycl@@AEBAXHP6A?AUkernel_param_desc_t@detail@23@H@Z@Z ?throw_asynchronous@queue@_V1@sycl@@QEAAXXZ ?unmap@experimental@oneapi@ext@_V1@sycl@@YAXPEBX_KAEBVcontext@45@@Z +?unmap_external_image_memory@experimental@oneapi@ext@_V1@sycl@@YAXUimage_mem_handle@12345@W4image_type@12345@AEBVdevice@45@AEBVcontext@45@@Z +?unmap_external_linear_memory@experimental@oneapi@ext@_V1@sycl@@YAXPEAXAEBVdevice@45@AEBVcontext@45@@Z ?unsampledImageConstructorNotification@detail@_V1@sycl@@YAXPEAX0AEBV?$optional@W4image_target@_V1@sycl@@@std@@W4mode@access@23@PEBXIAEBUcode_location@123@@Z ?unsampledImageConstructorNotification@image_plain@detail@_V1@sycl@@IEAAXAEBUcode_location@234@PEAXPEBXIQEA_KW4image_format@34@@Z ?unsampledImageDestructorNotification@image_plain@detail@_V1@sycl@@IEAAXPEAX@Z diff --git a/unified-runtime/include/ur_api.h b/unified-runtime/include/ur_api.h index e55b923663010..0d0396cc18374 100644 --- a/unified-runtime/include/ur_api.h +++ b/unified-runtime/include/ur_api.h @@ -467,6 +467,8 @@ typedef enum ur_function_t { UR_FUNCTION_BINDLESS_IMAGES_GET_IMAGE_SAMPLED_HANDLE_SUPPORT_EXP = 269, /// Enumerator for ::urBindlessImagesGetImageMemoryHandleTypeSupportExp UR_FUNCTION_BINDLESS_IMAGES_GET_IMAGE_MEMORY_HANDLE_TYPE_SUPPORT_EXP = 270, + /// Enumerator for ::urBindlessImagesFreeMappedLinearMemoryExp + UR_FUNCTION_BINDLESS_IMAGES_FREE_MAPPED_LINEAR_MEMORY_EXP = 271, /// @cond UR_FUNCTION_FORCE_UINT32 = 0x7fffffff /// @endcond @@ -10521,6 +10523,33 @@ UR_APIEXPORT ur_result_t UR_APICALL urBindlessImagesReleaseExternalMemoryExp( /// [in][release] handle of external memory to be destroyed ur_exp_external_mem_handle_t hExternalMem); +/////////////////////////////////////////////////////////////////////////////// +/// @brief Free a linear memory region mapped using MapExternalLinearMemoryExp +/// +/// @remarks +/// _Analogues_ +/// - **cuMemFree** +/// - **zeMemFree** +/// +/// @returns +/// - ::UR_RESULT_SUCCESS +/// - ::UR_RESULT_ERROR_UNINITIALIZED +/// - ::UR_RESULT_ERROR_DEVICE_LOST +/// - ::UR_RESULT_ERROR_ADAPTER_SPECIFIC +/// - ::UR_RESULT_ERROR_INVALID_NULL_HANDLE +/// + `NULL == hContext` +/// + `NULL == hDevice` +/// - ::UR_RESULT_ERROR_INVALID_CONTEXT +/// - ::UR_RESULT_ERROR_INVALID_NULL_POINTER +/// + `pMem == NULL` +UR_APIEXPORT ur_result_t UR_APICALL urBindlessImagesFreeMappedLinearMemoryExp( + /// [in] handle of the context object + ur_context_handle_t hContext, + /// [in] handle of the device object + ur_device_handle_t hDevice, + /// [in][release] pointer to mapped linear memory region to be freed + void *pMem); + /////////////////////////////////////////////////////////////////////////////// /// @brief Import an external semaphore /// @@ -14915,6 +14944,16 @@ typedef struct ur_bindless_images_release_external_memory_exp_params_t { ur_exp_external_mem_handle_t *phExternalMem; } ur_bindless_images_release_external_memory_exp_params_t; +/////////////////////////////////////////////////////////////////////////////// +/// @brief Function parameters for urBindlessImagesFreeMappedLinearMemoryExp +/// @details Each entry is a pointer to the parameter passed to the function; +/// allowing the callback the ability to modify the parameter's value +typedef struct ur_bindless_images_free_mapped_linear_memory_exp_params_t { + ur_context_handle_t *phContext; + ur_device_handle_t *phDevice; + void **ppMem; +} ur_bindless_images_free_mapped_linear_memory_exp_params_t; + /////////////////////////////////////////////////////////////////////////////// /// @brief Function parameters for urBindlessImagesImportExternalSemaphoreExp /// @details Each entry is a pointer to the parameter passed to the function; diff --git a/unified-runtime/include/ur_api_funcs.def b/unified-runtime/include/ur_api_funcs.def index 2ba60864f5940..af2c09896e18a 100644 --- a/unified-runtime/include/ur_api_funcs.def +++ b/unified-runtime/include/ur_api_funcs.def @@ -179,6 +179,7 @@ _UR_API(urBindlessImagesImportExternalMemoryExp) _UR_API(urBindlessImagesMapExternalArrayExp) _UR_API(urBindlessImagesMapExternalLinearMemoryExp) _UR_API(urBindlessImagesReleaseExternalMemoryExp) +_UR_API(urBindlessImagesFreeMappedLinearMemoryExp) _UR_API(urBindlessImagesImportExternalSemaphoreExp) _UR_API(urBindlessImagesReleaseExternalSemaphoreExp) _UR_API(urBindlessImagesWaitExternalSemaphoreExp) diff --git a/unified-runtime/include/ur_ddi.h b/unified-runtime/include/ur_ddi.h index 553bb61a7e8c7..6aefa464b78a4 100644 --- a/unified-runtime/include/ur_ddi.h +++ b/unified-runtime/include/ur_ddi.h @@ -1540,6 +1540,12 @@ typedef ur_result_t( typedef ur_result_t(UR_APICALL *ur_pfnBindlessImagesReleaseExternalMemoryExp_t)( ur_context_handle_t, ur_device_handle_t, ur_exp_external_mem_handle_t); +/////////////////////////////////////////////////////////////////////////////// +/// @brief Function-pointer for urBindlessImagesFreeMappedLinearMemoryExp +typedef ur_result_t( + UR_APICALL *ur_pfnBindlessImagesFreeMappedLinearMemoryExp_t)( + ur_context_handle_t, ur_device_handle_t, void *); + /////////////////////////////////////////////////////////////////////////////// /// @brief Function-pointer for urBindlessImagesImportExternalSemaphoreExp typedef ur_result_t( @@ -1593,6 +1599,7 @@ typedef struct ur_bindless_images_exp_dditable_t { ur_pfnBindlessImagesMapExternalLinearMemoryExp_t pfnMapExternalLinearMemoryExp; ur_pfnBindlessImagesReleaseExternalMemoryExp_t pfnReleaseExternalMemoryExp; + ur_pfnBindlessImagesFreeMappedLinearMemoryExp_t pfnFreeMappedLinearMemoryExp; ur_pfnBindlessImagesImportExternalSemaphoreExp_t pfnImportExternalSemaphoreExp; ur_pfnBindlessImagesReleaseExternalSemaphoreExp_t diff --git a/unified-runtime/include/ur_print.h b/unified-runtime/include/ur_print.h index c2dd79cc6ba23..92fd7a76a9a89 100644 --- a/unified-runtime/include/ur_print.h +++ b/unified-runtime/include/ur_print.h @@ -3183,6 +3183,19 @@ urPrintBindlessImagesReleaseExternalMemoryExpParams( *params, char *buffer, const size_t buff_size, size_t *out_size); +/////////////////////////////////////////////////////////////////////////////// +/// @brief Print ur_bindless_images_free_mapped_linear_memory_exp_params_t +/// struct +/// @returns +/// - ::UR_RESULT_SUCCESS +/// - ::UR_RESULT_ERROR_INVALID_SIZE +/// - `buff_size < out_size` +UR_APIEXPORT ur_result_t UR_APICALL +urPrintBindlessImagesFreeMappedLinearMemoryExpParams( + const struct ur_bindless_images_free_mapped_linear_memory_exp_params_t + *params, + char *buffer, const size_t buff_size, size_t *out_size); + /////////////////////////////////////////////////////////////////////////////// /// @brief Print ur_bindless_images_import_external_semaphore_exp_params_t /// struct diff --git a/unified-runtime/include/ur_print.hpp b/unified-runtime/include/ur_print.hpp index f73d2bb77cd71..f5e5b89dbb8ae 100644 --- a/unified-runtime/include/ur_print.hpp +++ b/unified-runtime/include/ur_print.hpp @@ -1256,6 +1256,9 @@ inline std::ostream &operator<<(std::ostream &os, enum ur_function_t value) { os << "UR_FUNCTION_BINDLESS_IMAGES_GET_IMAGE_MEMORY_HANDLE_TYPE_SUPPORT_" "EXP"; break; + case UR_FUNCTION_BINDLESS_IMAGES_FREE_MAPPED_LINEAR_MEMORY_EXP: + os << "UR_FUNCTION_BINDLESS_IMAGES_FREE_MAPPED_LINEAR_MEMORY_EXP"; + break; default: os << "unknown enumerator"; break; @@ -18627,6 +18630,32 @@ operator<<(std::ostream &os, [[maybe_unused]] const struct return os; } +/////////////////////////////////////////////////////////////////////////////// +/// @brief Print operator for the +/// ur_bindless_images_free_mapped_linear_memory_exp_params_t type +/// @returns +/// std::ostream & +inline std::ostream & +operator<<(std::ostream &os, [[maybe_unused]] const struct + ur_bindless_images_free_mapped_linear_memory_exp_params_t *params) { + + os << ".hContext = "; + + ur::details::printPtr(os, *(params->phContext)); + + os << ", "; + os << ".hDevice = "; + + ur::details::printPtr(os, *(params->phDevice)); + + os << ", "; + os << ".pMem = "; + + os << *(params->ppMem); + + return os; +} + /////////////////////////////////////////////////////////////////////////////// /// @brief Print operator for the /// ur_bindless_images_import_external_semaphore_exp_params_t type @@ -21414,6 +21443,10 @@ inline ur_result_t UR_APICALL printFunctionParams(std::ostream &os, os << (const struct ur_bindless_images_release_external_memory_exp_params_t *)params; } break; + case UR_FUNCTION_BINDLESS_IMAGES_FREE_MAPPED_LINEAR_MEMORY_EXP: { + os << (const struct + ur_bindless_images_free_mapped_linear_memory_exp_params_t *)params; + } break; case UR_FUNCTION_BINDLESS_IMAGES_IMPORT_EXTERNAL_SEMAPHORE_EXP: { os << (const struct ur_bindless_images_import_external_semaphore_exp_params_t *)params; diff --git a/unified-runtime/scripts/core/EXP-BINDLESS-IMAGES.rst b/unified-runtime/scripts/core/EXP-BINDLESS-IMAGES.rst index 479afb87e4f28..9b2ece6ed291c 100644 --- a/unified-runtime/scripts/core/EXP-BINDLESS-IMAGES.rst +++ b/unified-runtime/scripts/core/EXP-BINDLESS-IMAGES.rst @@ -142,6 +142,7 @@ Enums * ${X}_FUNCTION_BINDLESS_IMAGES_IMPORT_EXTERNAL_MEMORY_EXP * ${X}_FUNCTION_BINDLESS_IMAGES_MAP_EXTERNAL_ARRAY_EXP * ${X}_FUNCTION_BINDLESS_IMAGES_RELEASE_EXTERNAL_MEMORY_EXP + * ${X}_FUNCTION_BINDLESS_IMAGES_FREE_MAPPED_LINEAR_MEMORY_EXP * ${X}_FUNCTION_BINDLESS_IMAGES_IMPORT_EXTERNAL_SEMAPHORE_EXP * ${X}_FUNCTION_BINDLESS_IMAGES_RELEASE_EXTERNAL_SEMAPHORE_EXP * ${X}_FUNCTION_BINDLESS_IMAGES_WAIT_EXTERNAL_SEMAPHORE_EXP @@ -198,6 +199,7 @@ Functions * ${x}BindlessImagesMapExternalArrayExp * ${x}BindlessImagesMapExternalLinearMemoryExp * ${x}BindlessImagesReleaseExternalMemoryExp + * ${x}BindlessImagesFreeMappedLinearMemoryExp * ${x}BindlessImagesImportExternalSemaphoreExp * ${x}BindlessImagesReleaseExternalSemaphoreExp * ${x}BindlessImagesWaitExternalSemaphoreExp @@ -286,6 +288,8 @@ Changelog | || - GetImageUnsampledHandleSupportExp | | || - GetImageSampledHandleSupportExp | +----------+-------------------------------------------------------------+ +| 23.0 | Added BindlessImagesFreeMappedLinearMemory function. | ++----------+-------------------------------------------------------------+ Contributors -------------------------------------------------------------------------------- diff --git a/unified-runtime/scripts/core/exp-bindless-images.yml b/unified-runtime/scripts/core/exp-bindless-images.yml index 51e9cacc87db0..d3a491c344bea 100644 --- a/unified-runtime/scripts/core/exp-bindless-images.yml +++ b/unified-runtime/scripts/core/exp-bindless-images.yml @@ -901,6 +901,29 @@ returns: - $X_RESULT_ERROR_INVALID_VALUE --- #-------------------------------------------------------------------------- type: function +desc: "Free a linear memory region mapped using MapExternalLinearMemoryExp" +class: $xBindlessImages +name: FreeMappedLinearMemoryExp +ordinal: "0" +analogue: + - "**cuMemFree**" + - "**zeMemFree**" +params: + - type: $x_context_handle_t + name: hContext + desc: "[in] handle of the context object" + - type: $x_device_handle_t + name: hDevice + desc: "[in] handle of the device object" + - type: void * + name: pMem + desc: "[in][release] pointer to mapped linear memory region to be freed" +returns: + - $X_RESULT_ERROR_INVALID_CONTEXT + - $X_RESULT_ERROR_INVALID_NULL_POINTER: + - "`pMem == NULL`" +--- #-------------------------------------------------------------------------- +type: function desc: "Import an external semaphore" class: $xBindlessImages name: ImportExternalSemaphoreExp diff --git a/unified-runtime/scripts/core/registry.yml b/unified-runtime/scripts/core/registry.yml index e9d030ddeb994..ee1cdefcc13d2 100644 --- a/unified-runtime/scripts/core/registry.yml +++ b/unified-runtime/scripts/core/registry.yml @@ -658,6 +658,9 @@ etors: - name: BINDLESS_IMAGES_GET_IMAGE_MEMORY_HANDLE_TYPE_SUPPORT_EXP desc: Enumerator for $xBindlessImagesGetImageMemoryHandleTypeSupportExp value: '270' +- name: BINDLESS_IMAGES_FREE_MAPPED_LINEAR_MEMORY_EXP + desc: Enumerator for $xBindlessImagesFreeMappedLinearMemoryExp + value: '271' --- type: enum desc: Defines structure types diff --git a/unified-runtime/source/adapters/cuda/image.cpp b/unified-runtime/source/adapters/cuda/image.cpp index c22ef8a2a7f14..a843ab164f66b 100644 --- a/unified-runtime/source/adapters/cuda/image.cpp +++ b/unified-runtime/source/adapters/cuda/image.cpp @@ -1605,6 +1605,25 @@ UR_APIEXPORT ur_result_t UR_APICALL urBindlessImagesReleaseExternalMemoryExp( return UR_RESULT_SUCCESS; } +UR_APIEXPORT ur_result_t UR_APICALL urBindlessImagesFreeMappedLinearMemoryExp( + ur_context_handle_t hContext, ur_device_handle_t hDevice, void *pMem) { + UR_ASSERT(std::find(hContext->getDevices().begin(), + hContext->getDevices().end(), + hDevice) != hContext->getDevices().end(), + UR_RESULT_ERROR_INVALID_CONTEXT); + UR_ASSERT(pMem, UR_RESULT_ERROR_INVALID_NULL_POINTER); + + try { + ScopedContext Active(hDevice); + UR_CHECK_ERROR(cuMemFree(reinterpret_cast(pMem))); + } catch (ur_result_t Err) { + return Err; + } catch (...) { + return UR_RESULT_ERROR_UNKNOWN; + } + return UR_RESULT_SUCCESS; +} + UR_APIEXPORT ur_result_t UR_APICALL urBindlessImagesImportExternalSemaphoreExp( ur_context_handle_t hContext, ur_device_handle_t hDevice, ur_exp_external_semaphore_type_t semHandleType, diff --git a/unified-runtime/source/adapters/cuda/ur_interface_loader.cpp b/unified-runtime/source/adapters/cuda/ur_interface_loader.cpp index 5159915f861a6..dd33fea9fce26 100644 --- a/unified-runtime/source/adapters/cuda/ur_interface_loader.cpp +++ b/unified-runtime/source/adapters/cuda/ur_interface_loader.cpp @@ -344,6 +344,8 @@ UR_DLLEXPORT ur_result_t UR_APICALL urGetBindlessImagesExpProcAddrTable( urBindlessImagesMapExternalLinearMemoryExp; pDdiTable->pfnReleaseExternalMemoryExp = urBindlessImagesReleaseExternalMemoryExp; + pDdiTable->pfnFreeMappedLinearMemoryExp = + urBindlessImagesFreeMappedLinearMemoryExp; pDdiTable->pfnImportExternalSemaphoreExp = urBindlessImagesImportExternalSemaphoreExp; pDdiTable->pfnReleaseExternalSemaphoreExp = diff --git a/unified-runtime/source/adapters/hip/image.cpp b/unified-runtime/source/adapters/hip/image.cpp index 12a261fb51834..e0152fee14206 100644 --- a/unified-runtime/source/adapters/hip/image.cpp +++ b/unified-runtime/source/adapters/hip/image.cpp @@ -1430,6 +1430,25 @@ UR_APIEXPORT ur_result_t UR_APICALL urBindlessImagesReleaseExternalMemoryExp( return UR_RESULT_SUCCESS; } +UR_APIEXPORT ur_result_t UR_APICALL urBindlessImagesFreeMappedLinearMemoryExp( + ur_context_handle_t hContext, ur_device_handle_t hDevice, void *pMem) { + UR_ASSERT(std::find(hContext->getDevices().begin(), + hContext->getDevices().end(), + hDevice) != hContext->getDevices().end(), + UR_RESULT_ERROR_INVALID_CONTEXT); + UR_ASSERT(pMem, UR_RESULT_ERROR_INVALID_NULL_POINTER); + + try { + ScopedDevice Active(hDevice); + UR_CHECK_ERROR(hipFree(static_cast(pMem))); + } catch (ur_result_t Err) { + return Err; + } catch (...) { + return UR_RESULT_ERROR_UNKNOWN; + } + return UR_RESULT_SUCCESS; +} + UR_APIEXPORT ur_result_t UR_APICALL urBindlessImagesImportExternalSemaphoreExp( ur_context_handle_t hContext, ur_device_handle_t hDevice, ur_exp_external_semaphore_type_t semHandleType, diff --git a/unified-runtime/source/adapters/hip/ur_interface_loader.cpp b/unified-runtime/source/adapters/hip/ur_interface_loader.cpp index b32dcff5ba89a..104815fbcdd33 100644 --- a/unified-runtime/source/adapters/hip/ur_interface_loader.cpp +++ b/unified-runtime/source/adapters/hip/ur_interface_loader.cpp @@ -341,6 +341,8 @@ UR_DLLEXPORT ur_result_t UR_APICALL urGetBindlessImagesExpProcAddrTable( urBindlessImagesMapExternalLinearMemoryExp; pDdiTable->pfnReleaseExternalMemoryExp = urBindlessImagesReleaseExternalMemoryExp; + pDdiTable->pfnFreeMappedLinearMemoryExp = + urBindlessImagesFreeMappedLinearMemoryExp; pDdiTable->pfnImportExternalSemaphoreExp = urBindlessImagesImportExternalSemaphoreExp; pDdiTable->pfnReleaseExternalSemaphoreExp = diff --git a/unified-runtime/source/adapters/level_zero/CMakeLists.txt b/unified-runtime/source/adapters/level_zero/CMakeLists.txt index aa0da8da6fcba..1f8be4c625ee9 100644 --- a/unified-runtime/source/adapters/level_zero/CMakeLists.txt +++ b/unified-runtime/source/adapters/level_zero/CMakeLists.txt @@ -181,7 +181,6 @@ if(UR_BUILD_ADAPTER_L0_V2) ${CMAKE_CURRENT_SOURCE_DIR}/v2/event_provider_counter.cpp ${CMAKE_CURRENT_SOURCE_DIR}/v2/event_provider_normal.cpp ${CMAKE_CURRENT_SOURCE_DIR}/v2/event.cpp - ${CMAKE_CURRENT_SOURCE_DIR}/v2/image.cpp ${CMAKE_CURRENT_SOURCE_DIR}/v2/kernel.cpp ${CMAKE_CURRENT_SOURCE_DIR}/v2/memory.cpp ${CMAKE_CURRENT_SOURCE_DIR}/v2/queue_api.cpp diff --git a/unified-runtime/source/adapters/level_zero/device.hpp b/unified-runtime/source/adapters/level_zero/device.hpp index 3ce9b48bb258e..1ca19ed80cd4f 100644 --- a/unified-runtime/source/adapters/level_zero/device.hpp +++ b/unified-runtime/source/adapters/level_zero/device.hpp @@ -51,7 +51,6 @@ enum ur_ze_external_memory_desc_type { struct ur_ze_external_memory_data { void *importExtensionDesc; - ur_mem_handle_t urMemoryHandle; enum ur_ze_external_memory_desc_type type; size_t size; }; diff --git a/unified-runtime/source/adapters/level_zero/image.cpp b/unified-runtime/source/adapters/level_zero/image.cpp index bdd089acae4a5..414b9bba3d6d9 100644 --- a/unified-runtime/source/adapters/level_zero/image.cpp +++ b/unified-runtime/source/adapters/level_zero/image.cpp @@ -20,20 +20,8 @@ #include "loader/ze_loader.h" -namespace {} // namespace - namespace ur::level_zero { -ur_result_t -urBindlessImagesImageFreeExp(ur_context_handle_t /*hContext*/, - ur_device_handle_t /*hDevice*/, - ur_exp_image_mem_native_handle_t hImageMem) { - auto Native = reinterpret_cast(hImageMem); - delete Native; - - return UR_RESULT_SUCCESS; -} - ur_result_t urBindlessImagesImageCopyExp( ur_queue_handle_t hQueue, const void *pSrc, void *pDst, const ur_image_desc_t *pSrcImageDesc, const ur_image_desc_t *pDstImageDesc, diff --git a/unified-runtime/source/adapters/level_zero/image_common.cpp b/unified-runtime/source/adapters/level_zero/image_common.cpp index 0c8a8e40b3438..daa4cfef286d7 100644 --- a/unified-runtime/source/adapters/level_zero/image_common.cpp +++ b/unified-runtime/source/adapters/level_zero/image_common.cpp @@ -1103,6 +1103,16 @@ ur_result_t urBindlessImagesImageAllocateExp( return UR_RESULT_SUCCESS; } +ur_result_t +urBindlessImagesImageFreeExp([[maybe_unused]] ur_context_handle_t hContext, + [[maybe_unused]] ur_device_handle_t hDevice, + ur_exp_image_mem_native_handle_t hImageMem) { + ur_bindless_mem_handle_t *urImg = + reinterpret_cast(hImageMem); + delete urImg; + return UR_RESULT_SUCCESS; +} + ur_result_t urBindlessImagesUnsampledImageCreateExp( ur_context_handle_t hContext, ur_device_handle_t hDevice, ur_exp_image_mem_native_handle_t hImageMem, @@ -1307,8 +1317,6 @@ ur_result_t urBindlessImagesMapExternalArrayExp( UR_CALL(createUrImgFromZeImage(hContext->getZeHandle(), hDevice->ZeDevice, ZeImageDesc, phImageMem)); - externalMemoryData->urMemoryHandle = - reinterpret_cast(*phImageMem); return UR_RESULT_SUCCESS; } @@ -1322,8 +1330,6 @@ ur_result_t urBindlessImagesReleaseExternalMemoryExp( struct ur_ze_external_memory_data *externalMemoryData = reinterpret_cast(hExternalMem); - UR_CALL(ur::level_zero::urMemRelease(externalMemoryData->urMemoryHandle)); - switch (externalMemoryData->type) { case UR_ZE_EXTERNAL_OPAQUE_FD: delete (reinterpret_cast( @@ -1504,7 +1510,7 @@ ur_result_t urBindlessImagesMapExternalLinearMemoryExp( uint64_t size, ur_exp_external_mem_handle_t hExternalMem, void **phRetMem) { UR_ASSERT(hContext && hDevice && hExternalMem, UR_RESULT_ERROR_INVALID_NULL_HANDLE); - UR_ASSERT(offset && size, UR_RESULT_ERROR_INVALID_BUFFER_SIZE); + UR_ASSERT(size, UR_RESULT_ERROR_INVALID_BUFFER_SIZE); struct ur_ze_external_memory_data *externalMemoryData = reinterpret_cast(hExternalMem); @@ -1517,9 +1523,9 @@ ur_result_t urBindlessImagesMapExternalLinearMemoryExp( allocDesc.pNext = externalMemoryData->importExtensionDesc; void *mappedMemory; - ze_result_t zeResult = - zeMemAllocDevice(hContext->getZeHandle(), &allocDesc, size, 1, - hDevice->ZeDevice, &mappedMemory); + ze_result_t zeResult = ZE_CALL_NOCHECK( + zeMemAllocDevice, (hContext->getZeHandle(), &allocDesc, size, 0, + hDevice->ZeDevice, &mappedMemory)); if (zeResult != ZE_RESULT_SUCCESS) { return UR_RESULT_ERROR_OUT_OF_RESOURCES; } @@ -1527,15 +1533,23 @@ ur_result_t urBindlessImagesMapExternalLinearMemoryExp( zeResult = zeContextMakeMemoryResident(hContext->getZeHandle(), hDevice->ZeDevice, mappedMemory, size); if (zeResult != ZE_RESULT_SUCCESS) { - zeMemFree(hContext->getZeHandle(), mappedMemory); + ZE_CALL_NOCHECK(zeMemFree, (hContext->getZeHandle(), mappedMemory)); return UR_RESULT_ERROR_UNKNOWN; } + *phRetMem = reinterpret_cast( reinterpret_cast(mappedMemory) + offset); - externalMemoryData->urMemoryHandle = - reinterpret_cast(*phRetMem); + return UR_RESULT_SUCCESS; +} + +ur_result_t urBindlessImagesFreeMappedLinearMemoryExp( + ur_context_handle_t hContext, [[maybe_unused]] ur_device_handle_t hDevice, + void *pMem) { + UR_ASSERT(hContext, UR_RESULT_ERROR_INVALID_NULL_HANDLE); + UR_ASSERT(pMem, UR_RESULT_ERROR_INVALID_NULL_POINTER); + ZE2UR_CALL(zeMemFree, (hContext->getZeHandle(), pMem)); return UR_RESULT_SUCCESS; } diff --git a/unified-runtime/source/adapters/level_zero/ur_interface_loader.cpp b/unified-runtime/source/adapters/level_zero/ur_interface_loader.cpp index 908d20d6d9305..c80d2736fe7a7 100644 --- a/unified-runtime/source/adapters/level_zero/ur_interface_loader.cpp +++ b/unified-runtime/source/adapters/level_zero/ur_interface_loader.cpp @@ -100,6 +100,8 @@ UR_APIEXPORT ur_result_t UR_APICALL urGetBindlessImagesExpProcAddrTable( ur::level_zero::urBindlessImagesMapExternalLinearMemoryExp; pDdiTable->pfnReleaseExternalMemoryExp = ur::level_zero::urBindlessImagesReleaseExternalMemoryExp; + pDdiTable->pfnFreeMappedLinearMemoryExp = + ur::level_zero::urBindlessImagesFreeMappedLinearMemoryExp; pDdiTable->pfnImportExternalSemaphoreExp = ur::level_zero::urBindlessImagesImportExternalSemaphoreExp; pDdiTable->pfnReleaseExternalSemaphoreExp = diff --git a/unified-runtime/source/adapters/level_zero/ur_interface_loader.hpp b/unified-runtime/source/adapters/level_zero/ur_interface_loader.hpp index 0213e94dc8c84..9b94171759633 100644 --- a/unified-runtime/source/adapters/level_zero/ur_interface_loader.hpp +++ b/unified-runtime/source/adapters/level_zero/ur_interface_loader.hpp @@ -594,6 +594,8 @@ ur_result_t urBindlessImagesMapExternalLinearMemoryExp( ur_result_t urBindlessImagesReleaseExternalMemoryExp( ur_context_handle_t hContext, ur_device_handle_t hDevice, ur_exp_external_mem_handle_t hExternalMem); +ur_result_t urBindlessImagesFreeMappedLinearMemoryExp( + ur_context_handle_t hContext, ur_device_handle_t hDevice, void *pMem); ur_result_t urBindlessImagesImportExternalSemaphoreExp( ur_context_handle_t hContext, ur_device_handle_t hDevice, ur_exp_external_semaphore_type_t semHandleType, diff --git a/unified-runtime/source/adapters/level_zero/v2/image.cpp b/unified-runtime/source/adapters/level_zero/v2/image.cpp deleted file mode 100644 index 3a2af83458ff9..0000000000000 --- a/unified-runtime/source/adapters/level_zero/v2/image.cpp +++ /dev/null @@ -1,33 +0,0 @@ -//===--------- image.cpp - Level Zero Adapter -----------------------------===// -// -// Copyright (C) 2023 Intel Corporation -// -// Part of the Unified-Runtime Project, under the Apache License v2.0 with LLVM -// Exceptions. See LICENSE.TXT -// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception -// -//===----------------------------------------------------------------------===// - -#include "common.hpp" - -#include "../image_common.hpp" -#include "../ur_interface_loader.hpp" -#include "../v2/context.hpp" -#include "../v2/memory.hpp" -#include "logger/ur_logger.hpp" -#include "queue_api.hpp" -#include "queue_handle.hpp" - -namespace ur::level_zero { - -ur_result_t -urBindlessImagesImageFreeExp([[maybe_unused]] ur_context_handle_t hContext, - [[maybe_unused]] ur_device_handle_t hDevice, - ur_exp_image_mem_native_handle_t hImageMem) { - ur_bindless_mem_handle_t *urImg = - reinterpret_cast(hImageMem); - delete urImg; - return UR_RESULT_SUCCESS; -} - -} // namespace ur::level_zero diff --git a/unified-runtime/source/adapters/mock/ur_mockddi.cpp b/unified-runtime/source/adapters/mock/ur_mockddi.cpp index 6634cde2000ff..a7977e775a38d 100644 --- a/unified-runtime/source/adapters/mock/ur_mockddi.cpp +++ b/unified-runtime/source/adapters/mock/ur_mockddi.cpp @@ -9061,6 +9061,56 @@ __urdlllocal ur_result_t UR_APICALL urBindlessImagesReleaseExternalMemoryExp( return exceptionToResult(std::current_exception()); } +/////////////////////////////////////////////////////////////////////////////// +/// @brief Intercept function for urBindlessImagesFreeMappedLinearMemoryExp +__urdlllocal ur_result_t UR_APICALL urBindlessImagesFreeMappedLinearMemoryExp( + /// [in] handle of the context object + ur_context_handle_t hContext, + /// [in] handle of the device object + ur_device_handle_t hDevice, + /// [in][release] pointer to mapped linear memory region to be freed + void *pMem) try { + ur_result_t result = UR_RESULT_SUCCESS; + + ur_bindless_images_free_mapped_linear_memory_exp_params_t params = { + &hContext, &hDevice, &pMem}; + + auto beforeCallback = reinterpret_cast( + mock::getCallbacks().get_before_callback( + "urBindlessImagesFreeMappedLinearMemoryExp")); + if (beforeCallback) { + result = beforeCallback(¶ms); + if (result != UR_RESULT_SUCCESS) { + return result; + } + } + + auto replaceCallback = reinterpret_cast( + mock::getCallbacks().get_replace_callback( + "urBindlessImagesFreeMappedLinearMemoryExp")); + if (replaceCallback) { + result = replaceCallback(¶ms); + } else { + + result = UR_RESULT_SUCCESS; + } + + if (result != UR_RESULT_SUCCESS) { + return result; + } + + auto afterCallback = reinterpret_cast( + mock::getCallbacks().get_after_callback( + "urBindlessImagesFreeMappedLinearMemoryExp")); + if (afterCallback) { + return afterCallback(¶ms); + } + + return result; +} catch (...) { + return exceptionToResult(std::current_exception()); +} + /////////////////////////////////////////////////////////////////////////////// /// @brief Intercept function for urBindlessImagesImportExternalSemaphoreExp __urdlllocal ur_result_t UR_APICALL urBindlessImagesImportExternalSemaphoreExp( @@ -12055,6 +12105,9 @@ UR_DLLEXPORT ur_result_t UR_APICALL urGetBindlessImagesExpProcAddrTable( pDdiTable->pfnReleaseExternalMemoryExp = driver::urBindlessImagesReleaseExternalMemoryExp; + pDdiTable->pfnFreeMappedLinearMemoryExp = + driver::urBindlessImagesFreeMappedLinearMemoryExp; + pDdiTable->pfnImportExternalSemaphoreExp = driver::urBindlessImagesImportExternalSemaphoreExp; diff --git a/unified-runtime/source/adapters/native_cpu/image.cpp b/unified-runtime/source/adapters/native_cpu/image.cpp index 8ad2354927afc..def270ff931e7 100644 --- a/unified-runtime/source/adapters/native_cpu/image.cpp +++ b/unified-runtime/source/adapters/native_cpu/image.cpp @@ -181,6 +181,12 @@ UR_APIEXPORT ur_result_t UR_APICALL urBindlessImagesReleaseExternalMemoryExp( return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; } +UR_APIEXPORT ur_result_t UR_APICALL urBindlessImagesFreeMappedLinearMemoryExp( + [[maybe_unused]] ur_context_handle_t hContext, + [[maybe_unused]] ur_device_handle_t hDevice, [[maybe_unused]] void *pMem) { + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; +} + UR_APIEXPORT ur_result_t UR_APICALL urBindlessImagesImportExternalSemaphoreExp( [[maybe_unused]] ur_context_handle_t hContext, [[maybe_unused]] ur_device_handle_t hDevice, diff --git a/unified-runtime/source/adapters/native_cpu/ur_interface_loader.cpp b/unified-runtime/source/adapters/native_cpu/ur_interface_loader.cpp index fda01970ea959..924c5c3d0dd35 100644 --- a/unified-runtime/source/adapters/native_cpu/ur_interface_loader.cpp +++ b/unified-runtime/source/adapters/native_cpu/ur_interface_loader.cpp @@ -335,6 +335,8 @@ UR_DLLEXPORT ur_result_t UR_APICALL urGetBindlessImagesExpProcAddrTable( urBindlessImagesMapExternalLinearMemoryExp; pDdiTable->pfnReleaseExternalMemoryExp = urBindlessImagesReleaseExternalMemoryExp; + pDdiTable->pfnFreeMappedLinearMemoryExp = + urBindlessImagesFreeMappedLinearMemoryExp; pDdiTable->pfnImportExternalSemaphoreExp = urBindlessImagesImportExternalSemaphoreExp; pDdiTable->pfnReleaseExternalSemaphoreExp = diff --git a/unified-runtime/source/adapters/opencl/image.cpp b/unified-runtime/source/adapters/opencl/image.cpp index 3be2e3dfc4054..94f5084cc0164 100644 --- a/unified-runtime/source/adapters/opencl/image.cpp +++ b/unified-runtime/source/adapters/opencl/image.cpp @@ -181,6 +181,12 @@ UR_APIEXPORT ur_result_t UR_APICALL urBindlessImagesReleaseExternalMemoryExp( return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; } +UR_APIEXPORT ur_result_t UR_APICALL urBindlessImagesFreeMappedLinearMemoryExp( + [[maybe_unused]] ur_context_handle_t hContext, + [[maybe_unused]] ur_device_handle_t hDevice, [[maybe_unused]] void *pMem) { + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; +} + UR_APIEXPORT ur_result_t UR_APICALL urBindlessImagesImportExternalSemaphoreExp( [[maybe_unused]] ur_context_handle_t hContext, [[maybe_unused]] ur_device_handle_t hDevice, diff --git a/unified-runtime/source/adapters/opencl/ur_interface_loader.cpp b/unified-runtime/source/adapters/opencl/ur_interface_loader.cpp index f4cf124ac7f7d..1d74d09f8f5ae 100644 --- a/unified-runtime/source/adapters/opencl/ur_interface_loader.cpp +++ b/unified-runtime/source/adapters/opencl/ur_interface_loader.cpp @@ -354,6 +354,8 @@ UR_DLLEXPORT ur_result_t UR_APICALL urGetBindlessImagesExpProcAddrTable( urBindlessImagesMapExternalLinearMemoryExp; pDdiTable->pfnReleaseExternalMemoryExp = urBindlessImagesReleaseExternalMemoryExp; + pDdiTable->pfnFreeMappedLinearMemoryExp = + urBindlessImagesFreeMappedLinearMemoryExp; pDdiTable->pfnImportExternalSemaphoreExp = urBindlessImagesImportExternalSemaphoreExp; pDdiTable->pfnReleaseExternalSemaphoreExp = diff --git a/unified-runtime/source/loader/layers/tracing/ur_trcddi.cpp b/unified-runtime/source/loader/layers/tracing/ur_trcddi.cpp index ac4b357b7eb61..c7f7f26f1ece2 100644 --- a/unified-runtime/source/loader/layers/tracing/ur_trcddi.cpp +++ b/unified-runtime/source/loader/layers/tracing/ur_trcddi.cpp @@ -7617,6 +7617,49 @@ __urdlllocal ur_result_t UR_APICALL urBindlessImagesReleaseExternalMemoryExp( return result; } +/////////////////////////////////////////////////////////////////////////////// +/// @brief Intercept function for urBindlessImagesFreeMappedLinearMemoryExp +__urdlllocal ur_result_t UR_APICALL urBindlessImagesFreeMappedLinearMemoryExp( + /// [in] handle of the context object + ur_context_handle_t hContext, + /// [in] handle of the device object + ur_device_handle_t hDevice, + /// [in][release] pointer to mapped linear memory region to be freed + void *pMem) { + auto pfnFreeMappedLinearMemoryExp = + getContext()->urDdiTable.BindlessImagesExp.pfnFreeMappedLinearMemoryExp; + + if (nullptr == pfnFreeMappedLinearMemoryExp) + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; + + ur_bindless_images_free_mapped_linear_memory_exp_params_t params = { + &hContext, &hDevice, &pMem}; + uint64_t instance = getContext()->notify_begin( + UR_FUNCTION_BINDLESS_IMAGES_FREE_MAPPED_LINEAR_MEMORY_EXP, + "urBindlessImagesFreeMappedLinearMemoryExp", ¶ms); + + auto &logger = getContext()->logger; + UR_LOG_L(logger, INFO, " ---> urBindlessImagesFreeMappedLinearMemoryExp\n"); + + ur_result_t result = pfnFreeMappedLinearMemoryExp(hContext, hDevice, pMem); + + getContext()->notify_end( + UR_FUNCTION_BINDLESS_IMAGES_FREE_MAPPED_LINEAR_MEMORY_EXP, + "urBindlessImagesFreeMappedLinearMemoryExp", ¶ms, &result, instance); + + if (logger.getLevel() <= UR_LOGGER_LEVEL_INFO) { + std::ostringstream args_str; + ur::extras::printFunctionParams( + args_str, UR_FUNCTION_BINDLESS_IMAGES_FREE_MAPPED_LINEAR_MEMORY_EXP, + ¶ms); + UR_LOG_L(logger, INFO, + " <--- urBindlessImagesFreeMappedLinearMemoryExp({}) -> {};\n", + args_str.str(), result); + } + + return result; +} + /////////////////////////////////////////////////////////////////////////////// /// @brief Intercept function for urBindlessImagesImportExternalSemaphoreExp __urdlllocal ur_result_t UR_APICALL urBindlessImagesImportExternalSemaphoreExp( @@ -10268,6 +10311,11 @@ __urdlllocal ur_result_t UR_APICALL urGetBindlessImagesExpProcAddrTable( pDdiTable->pfnReleaseExternalMemoryExp = ur_tracing_layer::urBindlessImagesReleaseExternalMemoryExp; + dditable.pfnFreeMappedLinearMemoryExp = + pDdiTable->pfnFreeMappedLinearMemoryExp; + pDdiTable->pfnFreeMappedLinearMemoryExp = + ur_tracing_layer::urBindlessImagesFreeMappedLinearMemoryExp; + dditable.pfnImportExternalSemaphoreExp = pDdiTable->pfnImportExternalSemaphoreExp; pDdiTable->pfnImportExternalSemaphoreExp = diff --git a/unified-runtime/source/loader/layers/validation/ur_valddi.cpp b/unified-runtime/source/loader/layers/validation/ur_valddi.cpp index cfa5808f34944..d0c29b0987d77 100644 --- a/unified-runtime/source/loader/layers/validation/ur_valddi.cpp +++ b/unified-runtime/source/loader/layers/validation/ur_valddi.cpp @@ -8423,6 +8423,48 @@ __urdlllocal ur_result_t UR_APICALL urBindlessImagesReleaseExternalMemoryExp( return result; } +/////////////////////////////////////////////////////////////////////////////// +/// @brief Intercept function for urBindlessImagesFreeMappedLinearMemoryExp +__urdlllocal ur_result_t UR_APICALL urBindlessImagesFreeMappedLinearMemoryExp( + /// [in] handle of the context object + ur_context_handle_t hContext, + /// [in] handle of the device object + ur_device_handle_t hDevice, + /// [in][release] pointer to mapped linear memory region to be freed + void *pMem) { + auto pfnFreeMappedLinearMemoryExp = + getContext()->urDdiTable.BindlessImagesExp.pfnFreeMappedLinearMemoryExp; + + if (nullptr == pfnFreeMappedLinearMemoryExp) { + return UR_RESULT_ERROR_UNINITIALIZED; + } + + if (getContext()->enableParameterValidation) { + if (pMem == NULL) + return UR_RESULT_ERROR_INVALID_NULL_POINTER; + + if (NULL == hContext) + return UR_RESULT_ERROR_INVALID_NULL_HANDLE; + + if (NULL == hDevice) + return UR_RESULT_ERROR_INVALID_NULL_HANDLE; + } + + if (getContext()->enableLifetimeValidation && + !getContext()->refCountContext->isReferenceValid(hContext)) { + URLOG_CTX_INVALID_REFERENCE(hContext); + } + + if (getContext()->enableLifetimeValidation && + !getContext()->refCountContext->isReferenceValid(hDevice)) { + URLOG_CTX_INVALID_REFERENCE(hDevice); + } + + ur_result_t result = pfnFreeMappedLinearMemoryExp(hContext, hDevice, pMem); + + return result; +} + /////////////////////////////////////////////////////////////////////////////// /// @brief Intercept function for urBindlessImagesImportExternalSemaphoreExp __urdlllocal ur_result_t UR_APICALL urBindlessImagesImportExternalSemaphoreExp( @@ -11074,6 +11116,11 @@ UR_DLLEXPORT ur_result_t UR_APICALL urGetBindlessImagesExpProcAddrTable( pDdiTable->pfnReleaseExternalMemoryExp = ur_validation_layer::urBindlessImagesReleaseExternalMemoryExp; + dditable.pfnFreeMappedLinearMemoryExp = + pDdiTable->pfnFreeMappedLinearMemoryExp; + pDdiTable->pfnFreeMappedLinearMemoryExp = + ur_validation_layer::urBindlessImagesFreeMappedLinearMemoryExp; + dditable.pfnImportExternalSemaphoreExp = pDdiTable->pfnImportExternalSemaphoreExp; pDdiTable->pfnImportExternalSemaphoreExp = diff --git a/unified-runtime/source/loader/loader.def.in b/unified-runtime/source/loader/loader.def.in index 52b9fc18f863e..5ea8b003c1bc0 100644 --- a/unified-runtime/source/loader/loader.def.in +++ b/unified-runtime/source/loader/loader.def.in @@ -7,6 +7,7 @@ EXPORTS urAdapterRetain urAdapterSetLoggerCallback urAdapterSetLoggerCallbackLevel + urBindlessImagesFreeMappedLinearMemoryExp urBindlessImagesGetImageMemoryHandleTypeSupportExp urBindlessImagesGetImageSampledHandleSupportExp urBindlessImagesGetImageUnsampledHandleSupportExp @@ -190,6 +191,7 @@ EXPORTS urPrintApiVersion urPrintBaseDesc urPrintBaseProperties + urPrintBindlessImagesFreeMappedLinearMemoryExpParams urPrintBindlessImagesGetImageMemoryHandleTypeSupportExpParams urPrintBindlessImagesGetImageSampledHandleSupportExpParams urPrintBindlessImagesGetImageUnsampledHandleSupportExpParams diff --git a/unified-runtime/source/loader/loader.map.in b/unified-runtime/source/loader/loader.map.in index d1a7c8d190315..4838f1b6a12c7 100644 --- a/unified-runtime/source/loader/loader.map.in +++ b/unified-runtime/source/loader/loader.map.in @@ -7,6 +7,7 @@ urAdapterRetain; urAdapterSetLoggerCallback; urAdapterSetLoggerCallbackLevel; + urBindlessImagesFreeMappedLinearMemoryExp; urBindlessImagesGetImageMemoryHandleTypeSupportExp; urBindlessImagesGetImageSampledHandleSupportExp; urBindlessImagesGetImageUnsampledHandleSupportExp; @@ -190,6 +191,7 @@ urPrintApiVersion; urPrintBaseDesc; urPrintBaseProperties; + urPrintBindlessImagesFreeMappedLinearMemoryExpParams; urPrintBindlessImagesGetImageMemoryHandleTypeSupportExpParams; urPrintBindlessImagesGetImageSampledHandleSupportExpParams; urPrintBindlessImagesGetImageUnsampledHandleSupportExpParams; diff --git a/unified-runtime/source/loader/ur_ldrddi.cpp b/unified-runtime/source/loader/ur_ldrddi.cpp index 8ba14d7b8ad1c..8ef4d81d6e54b 100644 --- a/unified-runtime/source/loader/ur_ldrddi.cpp +++ b/unified-runtime/source/loader/ur_ldrddi.cpp @@ -7526,6 +7526,38 @@ __urdlllocal ur_result_t UR_APICALL urBindlessImagesReleaseExternalMemoryExp( return result; } +/////////////////////////////////////////////////////////////////////////////// +/// @brief Intercept function for urBindlessImagesFreeMappedLinearMemoryExp +__urdlllocal ur_result_t UR_APICALL urBindlessImagesFreeMappedLinearMemoryExp( + /// [in] handle of the context object + ur_context_handle_t hContext, + /// [in] handle of the device object + ur_device_handle_t hDevice, + /// [in][release] pointer to mapped linear memory region to be freed + void *pMem) { + ur_result_t result = UR_RESULT_SUCCESS; + + [[maybe_unused]] auto context = getContext(); + + // extract platform's function pointer table + auto dditable = reinterpret_cast(hContext)->dditable; + auto pfnFreeMappedLinearMemoryExp = + dditable->ur.BindlessImagesExp.pfnFreeMappedLinearMemoryExp; + if (nullptr == pfnFreeMappedLinearMemoryExp) + return UR_RESULT_ERROR_UNINITIALIZED; + + // convert loader handle to platform handle + hContext = reinterpret_cast(hContext)->handle; + + // convert loader handle to platform handle + hDevice = reinterpret_cast(hDevice)->handle; + + // forward to device-platform + result = pfnFreeMappedLinearMemoryExp(hContext, hDevice, pMem); + + return result; +} + /////////////////////////////////////////////////////////////////////////////// /// @brief Intercept function for urBindlessImagesImportExternalSemaphoreExp __urdlllocal ur_result_t UR_APICALL urBindlessImagesImportExternalSemaphoreExp( @@ -10294,6 +10326,8 @@ UR_DLLEXPORT ur_result_t UR_APICALL urGetBindlessImagesExpProcAddrTable( ur_loader::urBindlessImagesMapExternalLinearMemoryExp; pDdiTable->pfnReleaseExternalMemoryExp = ur_loader::urBindlessImagesReleaseExternalMemoryExp; + pDdiTable->pfnFreeMappedLinearMemoryExp = + ur_loader::urBindlessImagesFreeMappedLinearMemoryExp; pDdiTable->pfnImportExternalSemaphoreExp = ur_loader::urBindlessImagesImportExternalSemaphoreExp; pDdiTable->pfnReleaseExternalSemaphoreExp = diff --git a/unified-runtime/source/loader/ur_libapi.cpp b/unified-runtime/source/loader/ur_libapi.cpp index 7f07c0f4b21c7..3fcd7462206bd 100644 --- a/unified-runtime/source/loader/ur_libapi.cpp +++ b/unified-runtime/source/loader/ur_libapi.cpp @@ -8142,6 +8142,43 @@ ur_result_t UR_APICALL urBindlessImagesReleaseExternalMemoryExp( return exceptionToResult(std::current_exception()); } +/////////////////////////////////////////////////////////////////////////////// +/// @brief Free a linear memory region mapped using MapExternalLinearMemoryExp +/// +/// @remarks +/// _Analogues_ +/// - **cuMemFree** +/// - **zeMemFree** +/// +/// @returns +/// - ::UR_RESULT_SUCCESS +/// - ::UR_RESULT_ERROR_UNINITIALIZED +/// - ::UR_RESULT_ERROR_DEVICE_LOST +/// - ::UR_RESULT_ERROR_ADAPTER_SPECIFIC +/// - ::UR_RESULT_ERROR_INVALID_NULL_HANDLE +/// + `NULL == hContext` +/// + `NULL == hDevice` +/// - ::UR_RESULT_ERROR_INVALID_CONTEXT +/// - ::UR_RESULT_ERROR_INVALID_NULL_POINTER +/// + `pMem == NULL` +ur_result_t UR_APICALL urBindlessImagesFreeMappedLinearMemoryExp( + /// [in] handle of the context object + ur_context_handle_t hContext, + /// [in] handle of the device object + ur_device_handle_t hDevice, + /// [in][release] pointer to mapped linear memory region to be freed + void *pMem) try { + auto pfnFreeMappedLinearMemoryExp = + ur_lib::getContext() + ->urDdiTable.BindlessImagesExp.pfnFreeMappedLinearMemoryExp; + if (nullptr == pfnFreeMappedLinearMemoryExp) + return UR_RESULT_ERROR_UNINITIALIZED; + + return pfnFreeMappedLinearMemoryExp(hContext, hDevice, pMem); +} catch (...) { + return exceptionToResult(std::current_exception()); +} + /////////////////////////////////////////////////////////////////////////////// /// @brief Import an external semaphore /// diff --git a/unified-runtime/source/loader/ur_print.cpp b/unified-runtime/source/loader/ur_print.cpp index eabc4327735a5..3ea28e9fe9aeb 100644 --- a/unified-runtime/source/loader/ur_print.cpp +++ b/unified-runtime/source/loader/ur_print.cpp @@ -1373,6 +1373,15 @@ ur_result_t urPrintBindlessImagesReleaseExternalMemoryExpParams( return str_copy(&ss, buffer, buff_size, out_size); } +ur_result_t urPrintBindlessImagesFreeMappedLinearMemoryExpParams( + const struct ur_bindless_images_free_mapped_linear_memory_exp_params_t + *params, + char *buffer, const size_t buff_size, size_t *out_size) { + std::stringstream ss; + ss << params; + return str_copy(&ss, buffer, buff_size, out_size); +} + ur_result_t urPrintBindlessImagesImportExternalSemaphoreExpParams( const struct ur_bindless_images_import_external_semaphore_exp_params_t *params, diff --git a/unified-runtime/source/ur_api.cpp b/unified-runtime/source/ur_api.cpp index f0b300fa5f7b2..f11d9c4b6220c 100644 --- a/unified-runtime/source/ur_api.cpp +++ b/unified-runtime/source/ur_api.cpp @@ -7104,6 +7104,36 @@ ur_result_t UR_APICALL urBindlessImagesReleaseExternalMemoryExp( return result; } +/////////////////////////////////////////////////////////////////////////////// +/// @brief Free a linear memory region mapped using MapExternalLinearMemoryExp +/// +/// @remarks +/// _Analogues_ +/// - **cuMemFree** +/// - **zeMemFree** +/// +/// @returns +/// - ::UR_RESULT_SUCCESS +/// - ::UR_RESULT_ERROR_UNINITIALIZED +/// - ::UR_RESULT_ERROR_DEVICE_LOST +/// - ::UR_RESULT_ERROR_ADAPTER_SPECIFIC +/// - ::UR_RESULT_ERROR_INVALID_NULL_HANDLE +/// + `NULL == hContext` +/// + `NULL == hDevice` +/// - ::UR_RESULT_ERROR_INVALID_CONTEXT +/// - ::UR_RESULT_ERROR_INVALID_NULL_POINTER +/// + `pMem == NULL` +ur_result_t UR_APICALL urBindlessImagesFreeMappedLinearMemoryExp( + /// [in] handle of the context object + ur_context_handle_t hContext, + /// [in] handle of the device object + ur_device_handle_t hDevice, + /// [in][release] pointer to mapped linear memory region to be freed + void *pMem) { + ur_result_t result = UR_RESULT_SUCCESS; + return result; +} + /////////////////////////////////////////////////////////////////////////////// /// @brief Import an external semaphore ///