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<NDims> dims, sycl::range<NDims> 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<NDims> dims, sycl::range<NDims> 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..e14f5f65230f0 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<const char *> 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<const char *> 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<const char *> enabledDeviceExtensions(
+ std::begin(requiredExtensions), std::end(requiredExtensions));
+
const auto UUID = dev.get_info<sycl::ext::intel::info::device::uuid>();
// 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<VkExtensionProperties> supportedDeviceExtensions;
getSupportedDeviceExtensions(supportedDeviceExtensions, vk_physical_device);
const bool hasRequiredExtensions = std::all_of(
@@ -271,10 +302,29 @@ 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.
+ 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: "
<< devProps2.properties.deviceName << std::endl;
@@ -348,9 +398,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 +420,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 +437,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 +633,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 +655,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 +673,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 +773,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 +820,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 +874,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<ur::level_zero::ddi_getter>,
bool zeDriverImmediateCommandListAppendFound{false};
bool ZeDriverEuCountExtensionFound{false};
bool ZeCopyOffloadExtensionSupported{false};
+ bool ZeBindlessImagesExtensionSupported{false};
// Cache UR devices for reuse
std::vector<std::unique_ptr<ur_device_handle_t_>> URDevicesCache;