From f1658865ea084be52e172cc88140f9b1291264f7 Mon Sep 17 00:00:00 2001 From: Przemek Malon Date: Wed, 28 May 2025 14:36:50 +0100 Subject: [PATCH 1/2] [Bindless][UR][E2E] Fix/improve Vulkan tests. Fix aspect query. Some Vulkan interop tests did not correctly check for the `ext-oneapi-bindless-images` aspect. This lead to failures when devices (e.g. iGPU) supported memory import, but did not support Bindless Images. Adding the `// REQUIRES` lit directive fixed this issue for iGPU, however, it caused problems for Intel BMG machines, which do not pass the aspect query as it is currently implemented. To fix this, the query now returns Bindless Image support based on extensions returned from `zeDriverGetExtensionProperties`. The `env` vars set by LIT in the `buffer_usm.cpp` test are unnecessary and have been removed. This was causing a problem on iGPU where the interop would usually work, but the test was failing due to the `env` variables set for Bindless Images. A number of improvements have been made to `vulkan_common.hpp`. There is now a clearer distinction between required and optional instance and device extensions. Optional extensions are loaded only when supported. Fetching function pointers for optional extensions is now protected so that devices with no support do not try to attempt to retrieve pointers to those functions. The `VK_IMAGE_USAGE_STORAGE_BIT` was removed from image creation as it was not necessary to the functionality of the tests. The size of the depth texture created in `depth_format.cpp` was increased due to larger minimum memory allocation size requirements on iGPU. --- .../vulkan_interop/buffer_usm.cpp | 2 +- .../vulkan_interop/depth_format.cpp | 13 +- .../vulkan_interop/mipmaps.cpp | 4 +- .../vulkan_interop/sampled_images.cpp | 14 +- .../sampled_images_semaphore.cpp | 1 + .../vulkan_interop/unsampled_images.cpp | 1 + .../unsampled_images_semaphore.cpp | 1 + .../unsampled_images_timeline_semaphore.cpp | 1 + .../vulkan_interop/vulkan_common.hpp | 133 +++++++++++++----- .../source/adapters/level_zero/device.cpp | 9 +- .../source/adapters/level_zero/platform.cpp | 6 + .../source/adapters/level_zero/platform.hpp | 1 + 12 files changed, 131 insertions(+), 55 deletions(-) diff --git a/sycl/test-e2e/bindless_images/vulkan_interop/buffer_usm.cpp b/sycl/test-e2e/bindless_images/vulkan_interop/buffer_usm.cpp index 9d93397a6c94e..82ac7b27e8931 100644 --- a/sycl/test-e2e/bindless_images/vulkan_interop/buffer_usm.cpp +++ b/sycl/test-e2e/bindless_images/vulkan_interop/buffer_usm.cpp @@ -2,7 +2,7 @@ // 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 +// RUN: %{run} %t.out /** * This test does not use any image specific APIs. 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 4b78f18900f56..afd4960919f06 100644 --- a/sycl/test-e2e/bindless_images/vulkan_interop/depth_format.cpp +++ b/sycl/test-e2e/bindless_images/vulkan_interop/depth_format.cpp @@ -1,3 +1,4 @@ +// REQUIRES: aspect-ext_oneapi_bindless_images // REQUIRES: aspect-ext_oneapi_external_memory_import || (windows && level_zero && aspect-ext_oneapi_bindless_images) // REQUIRES: vulkan @@ -128,9 +129,8 @@ bool runTest(const sycl::device &syclDevice, sycl::range<2> dims, { vkInputImage = vkutil::createImage(imgType, imgInFormat, imgExtent, VK_IMAGE_USAGE_TRANSFER_SRC_BIT | - VK_IMAGE_USAGE_TRANSFER_DST_BIT | - VK_IMAGE_USAGE_STORAGE_BIT, - 1); + VK_IMAGE_USAGE_TRANSFER_DST_BIT, + 1 /*mipLevels*/); VkMemoryRequirements memRequirements; auto inputImageMemoryTypeIndex = vkutil::getImageMemoryTypeIndex( vkInputImage, VK_MEMORY_PROPERTY_DEVICE_LOCAL_BIT, memRequirements); @@ -141,9 +141,8 @@ bool runTest(const sycl::device &syclDevice, sycl::range<2> dims, vkOutputImage = vkutil::createImage(imgType, imgOutFormat, imgExtent, VK_IMAGE_USAGE_TRANSFER_SRC_BIT | - VK_IMAGE_USAGE_TRANSFER_DST_BIT | - VK_IMAGE_USAGE_STORAGE_BIT, - 1); + VK_IMAGE_USAGE_TRANSFER_DST_BIT, + 1 /*mipLevels*/); VkMemoryRequirements outputMemRequirements; auto outputImageMemoryTypeIndex = vkutil::getImageMemoryTypeIndex( vkOutputImage, VK_MEMORY_PROPERTY_DEVICE_LOCAL_BIT, @@ -373,7 +372,7 @@ int main() { return EXIT_FAILURE; } - auto testPassed = runTest(syclDevice, {16, 16}, {16, 16}); + auto testPassed = runTest(syclDevice, {128, 128}, {16, 16}); if (vkutil::cleanup() != VK_SUCCESS) { std::cerr << "Cleanup failed!\n"; diff --git a/sycl/test-e2e/bindless_images/vulkan_interop/mipmaps.cpp b/sycl/test-e2e/bindless_images/vulkan_interop/mipmaps.cpp index d9e72f104f6be..4594e4da59ae1 100644 --- a/sycl/test-e2e/bindless_images/vulkan_interop/mipmaps.cpp +++ b/sycl/test-e2e/bindless_images/vulkan_interop/mipmaps.cpp @@ -1,3 +1,4 @@ +// REQUIRES: aspect-ext_oneapi_bindless_images // REQUIRES: aspect-ext_oneapi_external_memory_import // REQUIRES: aspect-ext_oneapi_mipmap // REQUIRES: vulkan @@ -267,8 +268,7 @@ bool run_test(sycl::range dims, sycl::range localSize, // 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 | - VK_IMAGE_USAGE_STORAGE_BIT, + VK_IMAGE_USAGE_TRANSFER_DST_BIT, mipLevels); VkMemoryRequirements memRequirements; auto inputImageMemoryTypeIndex = vkutil::getImageMemoryTypeIndex( 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 9c13b5b445438..319fe8b2a4f3f 100644 --- a/sycl/test-e2e/bindless_images/vulkan_interop/sampled_images.cpp +++ b/sycl/test-e2e/bindless_images/vulkan_interop/sampled_images.cpp @@ -1,3 +1,4 @@ +// REQUIRES: aspect-ext_oneapi_bindless_images // REQUIRES: aspect-ext_oneapi_external_memory_import || (windows && level_zero && aspect-ext_oneapi_bindless_images) // REQUIRES: vulkan @@ -302,14 +303,13 @@ bool run_test(sycl::range dims, sycl::range localSize, 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 | - VK_IMAGE_USAGE_STORAGE_BIT, - 1 /*mipLevels*/ + auto inputImage = vkutil::createImage( + imgType, format, {width, height, depth}, + VK_IMAGE_USAGE_TRANSFER_SRC_BIT | VK_IMAGE_USAGE_TRANSFER_DST_BIT, + 1 /*mipLevels*/ #ifdef ENABLE_LINEAR_TILING - , - true /*linearTiling*/ + , + true /*linearTiling*/ #endif ); VkMemoryRequirements memRequirements; diff --git a/sycl/test-e2e/bindless_images/vulkan_interop/sampled_images_semaphore.cpp b/sycl/test-e2e/bindless_images/vulkan_interop/sampled_images_semaphore.cpp index 9b3409defe00f..321c9a489455c 100644 --- a/sycl/test-e2e/bindless_images/vulkan_interop/sampled_images_semaphore.cpp +++ b/sycl/test-e2e/bindless_images/vulkan_interop/sampled_images_semaphore.cpp @@ -1,3 +1,4 @@ +// REQUIRES: aspect-ext_oneapi_bindless_images // REQUIRES: target-nvidia || (windows && level_zero && aspect-ext_oneapi_bindless_images) // REQUIRES: vulkan 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 96e03e7730e42..fc1a6b90911b8 100644 --- a/sycl/test-e2e/bindless_images/vulkan_interop/unsampled_images.cpp +++ b/sycl/test-e2e/bindless_images/vulkan_interop/unsampled_images.cpp @@ -1,3 +1,4 @@ +// REQUIRES: aspect-ext_oneapi_bindless_images // REQUIRES: aspect-ext_oneapi_external_memory_import || (windows && level_zero && aspect-ext_oneapi_bindless_images) // REQUIRES: vulkan diff --git a/sycl/test-e2e/bindless_images/vulkan_interop/unsampled_images_semaphore.cpp b/sycl/test-e2e/bindless_images/vulkan_interop/unsampled_images_semaphore.cpp index 2a2c907e58e2e..53497f0b0af6a 100644 --- a/sycl/test-e2e/bindless_images/vulkan_interop/unsampled_images_semaphore.cpp +++ b/sycl/test-e2e/bindless_images/vulkan_interop/unsampled_images_semaphore.cpp @@ -1,3 +1,4 @@ +// REQUIRES: aspect-ext_oneapi_bindless_images // REQUIRES: aspect-ext_oneapi_external_semaphore_import // REQUIRES: aspect-ext_oneapi_external_memory_import // REQUIRES: vulkan 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 ac410be34df1a..fff6de90d09f0 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 @@ -1,3 +1,4 @@ +// REQUIRES: aspect-ext_oneapi_bindless_images // REQUIRES: aspect-ext_oneapi_external_memory_import // REQUIRES: aspect-ext_oneapi_external_semaphore_import // REQUIRES: vulkan 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 85cfa272ba0d3..b8f09dc1a21c8 100644 --- a/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_common.hpp +++ b/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_common.hpp @@ -61,8 +61,11 @@ static VkCommandPool vk_transferCmdPool; static VkCommandBuffer vk_computeCmdBuffer; static VkCommandBuffer vk_transferCmdBuffers[2]; +static bool supportsDedicatedAllocation = false; static bool requiresDedicatedAllocation = false; +static bool supportsExternalSemaphore = false; + // A static debug callback function that relays messages from the Vulkan // validation layer to the terminal. static VKAPI_ATTR VkBool32 VKAPI_CALL @@ -137,8 +140,11 @@ VkResult setupInstance() { std::vector requiredInstanceExtensions = { VK_EXT_DEBUG_UTILS_EXTENSION_NAME, VK_KHR_GET_PHYSICAL_DEVICE_PROPERTIES_2_EXTENSION_NAME, - VK_KHR_EXTERNAL_MEMORY_CAPABILITIES_EXTENSION_NAME, - VK_KHR_EXTERNAL_SEMAPHORE_CAPABILITIES_EXTENSION_NAME}; + VK_KHR_EXTERNAL_MEMORY_CAPABILITIES_EXTENSION_NAME}; + + std::vector optionalInstanceExtensions = { + VK_KHR_EXTERNAL_SEMAPHORE_CAPABILITIES_EXTENSION_NAME, + VK_KHR_DEDICATED_ALLOCATION_EXTENSION_NAME}; // Make sure that our required instance extensions are supported by the // running Vulkan instance. @@ -146,8 +152,23 @@ VkResult setupInstance() { std::string requiredExtension = requiredInstanceExtensions[i]; if (std::find(supportedInstanceExtensions.begin(), supportedInstanceExtensions.end(), - requiredExtension) == supportedInstanceExtensions.end()) + requiredExtension) == supportedInstanceExtensions.end()) { return VK_ERROR_EXTENSION_NOT_PRESENT; + } + } + + // Add any optional instance extensions that are supported by the + // running Vulkan instance. + for (int i = 0; i < optionalInstanceExtensions.size(); ++i) { + std::string optionalExtension = optionalInstanceExtensions[i]; + if (std::find(supportedInstanceExtensions.begin(), + supportedInstanceExtensions.end(), + optionalExtension) != supportedInstanceExtensions.end()) { + requiredInstanceExtensions.push_back(optionalInstanceExtensions[i]); + if (optionalExtension == VK_KHR_DEDICATED_ALLOCATION_EXTENSION_NAME) { + supportsDedicatedAllocation = true; + } + } } // Create the vulkan instance with our required extensions and layers. @@ -227,16 +248,25 @@ VkResult setupDevice(const sycl::device &dev) { static constexpr const char *requiredExtensions[] = { VK_KHR_GET_MEMORY_REQUIREMENTS_2_EXTENSION_NAME, VK_KHR_EXTERNAL_MEMORY_EXTENSION_NAME, - VK_KHR_EXTERNAL_SEMAPHORE_EXTENSION_NAME, #ifdef _WIN32 VK_KHR_EXTERNAL_MEMORY_WIN32_EXTENSION_NAME, +#else + VK_KHR_EXTERNAL_MEMORY_FD_EXTENSION_NAME, +#endif + }; + + static constexpr const char *optionalExtensions[] = { + VK_KHR_EXTERNAL_SEMAPHORE_EXTENSION_NAME, +#ifdef _WIN32 VK_KHR_EXTERNAL_SEMAPHORE_WIN32_EXTENSION_NAME, #else VK_KHR_EXTERNAL_SEMAPHORE_FD_EXTENSION_NAME, - VK_KHR_EXTERNAL_MEMORY_FD_EXTENSION_NAME, #endif }; + std::vector enabledDeviceExtensions( + std::begin(requiredExtensions), std::end(requiredExtensions)); + const auto UUID = dev.get_info(); // From all physical devices, find the first one with a matching UUID @@ -259,6 +289,7 @@ VkResult setupDevice(const sycl::device &dev) { continue; } + // Check if the device supports the required extensions. std::vector supportedDeviceExtensions; getSupportedDeviceExtensions(supportedDeviceExtensions, vk_physical_device); const bool hasRequiredExtensions = std::all_of( @@ -271,10 +302,31 @@ VkResult setupDevice(const sycl::device &dev) { }); return (it != std::end(supportedDeviceExtensions)); }); + // Skip this device if it does not support all required extensions. if (!hasRequiredExtensions) { continue; } + // Check if the device supports the optional extensions, if so add them to + // the list of enabled device extensions. + std::for_each(std::begin(optionalExtensions), std::end(optionalExtensions), + [&](const char *optionalExt) -> void { + auto it = + std::find_if(std::begin(supportedDeviceExtensions), + std::end(supportedDeviceExtensions), + [&](const VkExtensionProperties &ext) { + return (ext.extensionName == + std::string_view(optionalExt)); + }); + if (it != std::end(supportedDeviceExtensions)) { + enabledDeviceExtensions.push_back(optionalExt); + if (std::string_view(optionalExt) == + VK_KHR_EXTERNAL_SEMAPHORE_EXTENSION_NAME) { + supportsExternalSemaphore = true; + } + } + }); + foundDevice = true; std::cout << "Found suitable Vulkan device: " << devProps2.properties.deviceName << std::endl; @@ -348,9 +400,8 @@ VkResult setupDevice(const sycl::device &dev) { dci.pQueueCreateInfos = qcis.data(); dci.queueCreateInfoCount = qcis.size(); dci.pEnabledFeatures = &deviceFeatures; - dci.enabledExtensionCount = - sizeof(requiredExtensions) / sizeof(requiredExtensions[0]); - dci.ppEnabledExtensionNames = &requiredExtensions[0]; + dci.enabledExtensionCount = enabledDeviceExtensions.size(); + dci.ppEnabledExtensionNames = enabledDeviceExtensions.data(); VK_CHECK_CALL_RET( vkCreateDevice(vk_physical_device, &dci, nullptr, &vk_device)); @@ -371,13 +422,15 @@ VkResult setupDevice(const sycl::device &dev) { << "Could not get func pointer to \"vkGetMemoryWin32HandleKHR\"!\n"; return VK_ERROR_UNKNOWN; } - vk_getSemaphoreWin32HandleKHR = - (PFN_vkGetSemaphoreWin32HandleKHR)vkGetDeviceProcAddr( - vk_device, "vkGetSemaphoreWin32HandleKHR"); - if (!vk_getSemaphoreWin32HandleKHR) { - std::cerr - << "Could not get func pointer to \"vkGetSemaphoreWin32HandleKHR\"!\n"; - return VK_ERROR_UNKNOWN; + if (supportsExternalSemaphore) { + vk_getSemaphoreWin32HandleKHR = + (PFN_vkGetSemaphoreWin32HandleKHR)vkGetDeviceProcAddr( + vk_device, "vkGetSemaphoreWin32HandleKHR"); + if (!vk_getSemaphoreWin32HandleKHR) { + std::cerr << "Could not get func pointer to " + "\"vkGetSemaphoreWin32HandleKHR\"!\n"; + return VK_ERROR_UNKNOWN; + } } #else vk_getMemoryFdKHR = @@ -386,11 +439,13 @@ VkResult setupDevice(const sycl::device &dev) { std::cerr << "Could not get func pointer to \"vkGetMemoryFdKHR\"!\n"; return VK_ERROR_UNKNOWN; } - vk_getSemaphoreFdKHR = (PFN_vkGetSemaphoreFdKHR)vkGetDeviceProcAddr( - vk_device, "vkGetSemaphoreFdKHR"); - if (!vk_getSemaphoreFdKHR) { - std::cerr << "Could not get func pointer to \"vkGetSemaphoreFdKHR\"!\n"; - return VK_ERROR_UNKNOWN; + if (supportsExternalSemaphore) { + vk_getSemaphoreFdKHR = (PFN_vkGetSemaphoreFdKHR)vkGetDeviceProcAddr( + vk_device, "vkGetSemaphoreFdKHR"); + if (!vk_getSemaphoreFdKHR) { + std::cerr << "Could not get func pointer to \"vkGetSemaphoreFdKHR\"!\n"; + return VK_ERROR_UNKNOWN; + } } #endif @@ -580,10 +635,11 @@ VkDeviceMemory allocateDeviceMemory(size_t size, uint32_t memoryTypeIndex, #else emai.handleTypes = VK_EXTERNAL_MEMORY_HANDLE_TYPE_OPAQUE_FD_BIT; #endif - if (requiresDedicatedAllocation) + if (requiresDedicatedAllocation) { dedicatedInfo.pNext = &emai; - else + } else { mai.pNext = &emai; + } } VkDeviceMemory memory; @@ -601,12 +657,15 @@ property flags passed. */ uint32_t getImageMemoryTypeIndex(VkImage image, VkMemoryPropertyFlags flags, VkMemoryRequirements &memRequirements) { - VkMemoryDedicatedRequirements dedicatedRequirements{}; - dedicatedRequirements.sType = VK_STRUCTURE_TYPE_MEMORY_DEDICATED_REQUIREMENTS; - VkMemoryRequirements2 memoryRequirements2{}; memoryRequirements2.sType = VK_STRUCTURE_TYPE_MEMORY_REQUIREMENTS_2; - memoryRequirements2.pNext = &dedicatedRequirements; + + VkMemoryDedicatedRequirements dedicatedRequirements{}; + if (supportsDedicatedAllocation) { + dedicatedRequirements.sType = + VK_STRUCTURE_TYPE_MEMORY_DEDICATED_REQUIREMENTS; + memoryRequirements2.pNext = &dedicatedRequirements; + } VkImageMemoryRequirementsInfo2 imageRequirementsInfo{}; imageRequirementsInfo.sType = @@ -616,8 +675,9 @@ uint32_t getImageMemoryTypeIndex(VkImage image, VkMemoryPropertyFlags flags, vk_getImageMemoryRequirements2(vk_device, &imageRequirementsInfo, &memoryRequirements2); - if (dedicatedRequirements.requiresDedicatedAllocation) + if (dedicatedRequirements.requiresDedicatedAllocation) { requiresDedicatedAllocation = true; + } VkPhysicalDeviceMemoryProperties memProperties; vkGetPhysicalDeviceMemoryProperties(vk_physical_device, &memProperties); @@ -715,6 +775,11 @@ HANDLE getSemaphoreWin32Handle(VkSemaphore semaphore) { sghwi.semaphore = semaphore; sghwi.handleType = VK_EXTERNAL_SEMAPHORE_HANDLE_TYPE_OPAQUE_WIN32_BIT; + if (!supportsExternalSemaphore) { + std::cerr << "External semaphore support is not enabled!\n"; + return 0; + } + if (vk_getSemaphoreWin32HandleKHR != nullptr) { VK_CHECK_CALL(vk_getSemaphoreWin32HandleKHR(vk_device, &sghwi, &retHandle)); } else { @@ -757,6 +822,12 @@ int getSemaphoreOpaqueFD(VkSemaphore semaphore) { sgfi.handleType = VK_EXTERNAL_SEMAPHORE_HANDLE_TYPE_OPAQUE_FD_BIT; int fd = 0; + + if (!supportsExternalSemaphore) { + std::cerr << "External semaphore support is not enabled!\n"; + return 0; + } + if (vk_getSemaphoreFdKHR != nullptr) { VK_CHECK_CALL(vk_getSemaphoreFdKHR(vk_device, &sgfi, &fd)); } else { @@ -805,11 +876,9 @@ struct vulkan_image_test_resources_t { vulkan_image_test_resources_t(VkImageType imgType, VkFormat format, VkExtent3D ext, const size_t imageSizeBytes) { - vkImage = vkutil::createImage(imgType, format, ext, - VK_IMAGE_USAGE_TRANSFER_SRC_BIT | - VK_IMAGE_USAGE_TRANSFER_DST_BIT | - VK_IMAGE_USAGE_STORAGE_BIT, - 1); + vkImage = vkutil::createImage( + imgType, format, ext, + VK_IMAGE_USAGE_TRANSFER_SRC_BIT | VK_IMAGE_USAGE_TRANSFER_DST_BIT, 1); VkMemoryRequirements memRequirements; auto inputImageMemoryTypeIndex = vkutil::getImageMemoryTypeIndex( vkImage, VK_MEMORY_PROPERTY_DEVICE_LOCAL_BIT, memRequirements); diff --git a/unified-runtime/source/adapters/level_zero/device.cpp b/unified-runtime/source/adapters/level_zero/device.cpp index c372443d07ec2..1d135b9466eb1 100644 --- a/unified-runtime/source/adapters/level_zero/device.cpp +++ b/unified-runtime/source/adapters/level_zero/device.cpp @@ -1087,21 +1087,18 @@ ur_result_t urDeviceGetInfo( case UR_DEVICE_INFO_COMMAND_BUFFER_SUBGRAPH_SUPPORT_EXP: return ReturnValue(false); case UR_DEVICE_INFO_BINDLESS_IMAGES_SUPPORT_EXP: { - return ReturnValue(Device->isIntelDG2OrNewer() && - Device->ZeDeviceImageProperties->maxImageDims1D > 0 && - Device->ZeDeviceImageProperties->maxImageDims2D > 0 && - Device->ZeDeviceImageProperties->maxImageDims3D > 0); + return ReturnValue(Device->Platform->ZeBindlessImagesExtensionSupported); } case UR_DEVICE_INFO_BINDLESS_IMAGES_SHARED_USM_SUPPORT_EXP: { // On L0 bindless images can not be backed by shared (managed) USM. return ReturnValue(false); } case UR_DEVICE_INFO_BINDLESS_IMAGES_1D_USM_SUPPORT_EXP: { - return ReturnValue(Device->isIntelDG2OrNewer() && + return ReturnValue(Device->Platform->ZeBindlessImagesExtensionSupported && Device->ZeDeviceImageProperties->maxImageDims1D > 0); } case UR_DEVICE_INFO_BINDLESS_IMAGES_2D_USM_SUPPORT_EXP: { - return ReturnValue(Device->isIntelDG2OrNewer() && + return ReturnValue(Device->Platform->ZeBindlessImagesExtensionSupported && Device->ZeDeviceImageProperties->maxImageDims2D > 0); } case UR_DEVICE_INFO_IMAGE_PITCH_ALIGN_EXP: diff --git a/unified-runtime/source/adapters/level_zero/platform.cpp b/unified-runtime/source/adapters/level_zero/platform.cpp index 2602f75e30ef3..50a2ce1080964 100644 --- a/unified-runtime/source/adapters/level_zero/platform.cpp +++ b/unified-runtime/source/adapters/level_zero/platform.cpp @@ -294,6 +294,12 @@ ur_result_t ur_platform_handle_t_::initialize() { ZeCopyOffloadExtensionSupported = true; } } + if (strncmp(extension.name, ZE_BINDLESS_IMAGE_EXP_NAME, + strlen(ZE_BINDLESS_IMAGE_EXP_NAME) + 1) == 0) { + if (extension.version == ZE_BINDLESS_IMAGE_EXP_VERSION_1_0) { + ZeBindlessImagesExtensionSupported = true; + } + } zeDriverExtensionMap[extension.name] = extension.version; } diff --git a/unified-runtime/source/adapters/level_zero/platform.hpp b/unified-runtime/source/adapters/level_zero/platform.hpp index 62b36867304a0..192cd063c19e5 100644 --- a/unified-runtime/source/adapters/level_zero/platform.hpp +++ b/unified-runtime/source/adapters/level_zero/platform.hpp @@ -68,6 +68,7 @@ struct ur_platform_handle_t_ : ur::handle_base, bool zeDriverImmediateCommandListAppendFound{false}; bool ZeDriverEuCountExtensionFound{false}; bool ZeCopyOffloadExtensionSupported{false}; + bool ZeBindlessImagesExtensionSupported{false}; // Cache UR devices for reuse std::vector> URDevicesCache; From 1b04c69aa48b5aa12a9260d11a134af10c787c7f Mon Sep 17 00:00:00 2001 From: Przemek Malon Date: Thu, 29 May 2025 11:07:04 +0100 Subject: [PATCH 2/2] Address feedback --- .../vulkan_interop/vulkan_common.hpp | 32 +++++++++---------- 1 file changed, 15 insertions(+), 17 deletions(-) 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 b8f09dc1a21c8..e14f5f65230f0 100644 --- a/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_common.hpp +++ b/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_common.hpp @@ -309,23 +309,21 @@ VkResult setupDevice(const sycl::device &dev) { // Check if the device supports the optional extensions, if so add them to // the list of enabled device extensions. - std::for_each(std::begin(optionalExtensions), std::end(optionalExtensions), - [&](const char *optionalExt) -> void { - auto it = - std::find_if(std::begin(supportedDeviceExtensions), - std::end(supportedDeviceExtensions), - [&](const VkExtensionProperties &ext) { - return (ext.extensionName == - std::string_view(optionalExt)); - }); - if (it != std::end(supportedDeviceExtensions)) { - enabledDeviceExtensions.push_back(optionalExt); - if (std::string_view(optionalExt) == - VK_KHR_EXTERNAL_SEMAPHORE_EXTENSION_NAME) { - supportsExternalSemaphore = true; - } - } - }); + for (const char *optionalExt : optionalExtensions) { + auto it = std::find_if(std::begin(supportedDeviceExtensions), + std::end(supportedDeviceExtensions), + [&](const VkExtensionProperties &ext) { + return (ext.extensionName == + std::string_view(optionalExt)); + }); + if (it != std::end(supportedDeviceExtensions)) { + enabledDeviceExtensions.push_back(optionalExt); + if (std::string_view(optionalExt) == + VK_KHR_EXTERNAL_SEMAPHORE_EXTENSION_NAME) { + supportsExternalSemaphore = true; + } + } + } foundDevice = true; std::cout << "Found suitable Vulkan device: "