diff --git a/.github/CODEOWNERS b/.github/CODEOWNERS index fb5054dadbc52..aecd7b2d843a8 100644 --- a/.github/CODEOWNERS +++ b/.github/CODEOWNERS @@ -190,7 +190,6 @@ sycl/source/detail/bindless* @intel/bindless-images-reviewers sycl/source/detail/memory_export.cpp @intel/bindless-images-reviewers sycl/test/check_device_code/extensions/bindless_images.cpp @intel/bindless-images-reviewers sycl/test-e2e/bindless_images/ @intel/bindless-images-reviewers -sycl/test-e2e/CommonUtils/vulkan_common.hpp @intel/bindless-images-reviewers sycl/test-e2e/MemoryExport/ @intel/bindless-images-reviewers sycl/unittests/Extensions/BindlessImages/ @intel/bindless-images-reviewers diff --git a/sycl/test-e2e/CommonUtils/vulkan_common.hpp b/sycl/test-e2e/CommonUtils/vulkan_common.hpp deleted file mode 100644 index 0a4498e758aaf..0000000000000 --- a/sycl/test-e2e/CommonUtils/vulkan_common.hpp +++ /dev/null @@ -1,1123 +0,0 @@ -#pragma once - -#ifdef _WIN32 -#define VK_USE_PLATFORM_WIN32_KHR -#endif - -#include - -#include -#include -#include -#include -#include -#include -#include -#include - -void printString(std::string str) { -#ifdef VERBOSE_PRINT - std::cout << str << std::endl; -#endif -} - -#define VK_CHECK_CALL_RET(call) \ - { \ - VkResult err = call; \ - if (err != VK_SUCCESS) { \ - std::cerr << #call << " failed. Code: " << err << "\n"; \ - return err; \ - } \ - } - -#define VK_CHECK_CALL(call) \ - { \ - VkResult err = call; \ - if (err != VK_SUCCESS) \ - std::cerr << #call << " failed. Code: " << err << "\n"; \ - } - -static VkInstance vk_instance; -static VkPhysicalDevice vk_physical_device; -static VkDebugUtilsMessengerEXT vk_debug_messenger; -static VkDevice vk_device; -static VkQueue vk_compute_queue; -static VkQueue vk_transfer_queue; - -#ifdef _WIN32 -static PFN_vkGetMemoryWin32HandleKHR vk_GetMemoryWin32HandleKHR; -static PFN_vkGetSemaphoreWin32HandleKHR vk_getSemaphoreWin32HandleKHR; -#else -static PFN_vkGetMemoryFdKHR vk_getMemoryFdKHR; -static PFN_vkGetSemaphoreFdKHR vk_getSemaphoreFdKHR; -#endif - -static PFN_vkGetImageMemoryRequirements2 vk_getImageMemoryRequirements2; - -static uint32_t vk_computeQueueFamilyIndex; -static uint32_t vk_transferQueueFamilyIndex; - -static VkCommandPool vk_computeCmdPool; -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; -static bool supportsDmaBuf = false; - -// A static debug callback function that relays messages from the Vulkan -// validation layer to the terminal. -static VKAPI_ATTR VkBool32 VKAPI_CALL -debugCallback(VkDebugUtilsMessageSeverityFlagBitsEXT messageSeverity, - VkDebugUtilsMessageTypeFlagsEXT messageType, - const VkDebugUtilsMessengerCallbackDataEXT *pCallbackData, - void *pUserData) { - // Only print errors from validation layer - if (messageSeverity & VK_DEBUG_UTILS_MESSAGE_SEVERITY_ERROR_BIT_EXT) { - std::cerr << pCallbackData->pMessage << "\n\n"; - } - return VK_FALSE; -} - -namespace vkutil { - -// Returns all supported Vulkan instance extensions. -VkResult -getSupportedInstanceExtensions(std::vector &supportedExtensions) { - uint32_t count = 0; - VK_CHECK_CALL_RET( - vkEnumerateInstanceExtensionProperties(nullptr, &count, nullptr)); - - std::vector extensionProperties(count); - - VK_CHECK_CALL_RET(vkEnumerateInstanceExtensionProperties( - nullptr, &count, extensionProperties.data())); - - for (auto &extension : extensionProperties) { - supportedExtensions.push_back(extension.extensionName); - } - - return VK_SUCCESS; -} - -/* -In this function we set up the Vulkan instance, which is the one of the first -steps in setting up a Vulkan application. -When creating an instance we need to specify some information about our -application, most importantly, we need to specify some extensions that we -require to perform interop operations. -*/ -VkResult setupInstance() { - // Generic application information. The specific values are not important to - // the execution of the Vulkan program. - VkApplicationInfo ai = {}; - ai.sType = VK_STRUCTURE_TYPE_APPLICATION_INFO; - ai.pApplicationName = "SYCL-Vulkan-Interop"; - ai.applicationVersion = VK_MAKE_VERSION(1, 0, 0); - ai.pEngineName = ""; - ai.engineVersion = VK_MAKE_VERSION(1, 0, 0); - ai.apiVersion = VK_API_VERSION_1_0; - - // Query the number of available layers and retrieve their names. One example - // of a layer is the validation layer, this layer allows for runtime debug - // messages to be returned if anything goes wrong in the Vulkan application. - // We will set up a callback function to print debug information if the - // validation layer is available. - uint32_t layerCount; - VK_CHECK_CALL_RET(vkEnumerateInstanceLayerProperties(&layerCount, nullptr)); - - std::vector availableLayers(layerCount); - VK_CHECK_CALL_RET( - vkEnumerateInstanceLayerProperties(&layerCount, availableLayers.data())); - - // Query the supported instance extensions. - std::vector supportedInstanceExtensions; - VK_CHECK_CALL_RET( - getSupportedInstanceExtensions(supportedInstanceExtensions)); - - // We have some instance extensions that we require for the tests to function. - 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}; - - 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. - for (int i = 0; i < requiredInstanceExtensions.size(); ++i) { - std::string requiredExtension = requiredInstanceExtensions[i]; - if (std::find(supportedInstanceExtensions.begin(), - 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. - VkInstanceCreateInfo ci = {}; - ci.sType = VK_STRUCTURE_TYPE_INSTANCE_CREATE_INFO; - ci.pApplicationInfo = &ai; - ci.enabledExtensionCount = requiredInstanceExtensions.size(); - ci.ppEnabledExtensionNames = requiredInstanceExtensions.data(); - const char *validationLayerName = "VK_LAYER_KHRONOS_validation"; - bool validationLayerAvailable = std::any_of( - availableLayers.begin(), availableLayers.end(), [&](const auto &layer) { - return strcmp(layer.layerName, validationLayerName) == 0; - }); - - if (validationLayerAvailable) { - ci.enabledLayerCount = 1; - ci.ppEnabledLayerNames = &validationLayerName; - } else { - std::cerr - << "VK_LAYER_KHRONOS_validation not present, validation is disabled \n"; - } - - VK_CHECK_CALL_RET(vkCreateInstance(&ci, nullptr, &vk_instance)); - - // Create a debug utils messenger. This will allow us to print debug - // information from the Vulkan validation layer. - VkDebugUtilsMessengerCreateInfoEXT dumci = {}; - dumci.sType = VK_STRUCTURE_TYPE_DEBUG_UTILS_MESSENGER_CREATE_INFO_EXT; - dumci.messageSeverity = VK_DEBUG_UTILS_MESSAGE_SEVERITY_VERBOSE_BIT_EXT | - VK_DEBUG_UTILS_MESSAGE_SEVERITY_WARNING_BIT_EXT | - VK_DEBUG_UTILS_MESSAGE_SEVERITY_ERROR_BIT_EXT; - dumci.messageType = VK_DEBUG_UTILS_MESSAGE_TYPE_GENERAL_BIT_EXT | - VK_DEBUG_UTILS_MESSAGE_TYPE_VALIDATION_BIT_EXT | - VK_DEBUG_UTILS_MESSAGE_TYPE_PERFORMANCE_BIT_EXT; - dumci.pfnUserCallback = debugCallback; - - auto func = (PFN_vkCreateDebugUtilsMessengerEXT)vkGetInstanceProcAddr( - vk_instance, "vkCreateDebugUtilsMessengerEXT"); - if (func != nullptr) { - VK_CHECK_CALL_RET(func(vk_instance, &dumci, nullptr, &vk_debug_messenger)); - } else { - return VK_ERROR_EXTENSION_NOT_PRESENT; - } - - return VK_SUCCESS; -} - -// Returns all supported Vulkan device extensions. -VkResult -getSupportedDeviceExtensions(std::vector &extensions, - VkPhysicalDevice device) { - uint32_t numExtensions = 0; - - VK_CHECK_CALL_RET(vkEnumerateDeviceExtensionProperties( - device, nullptr, &numExtensions, nullptr)); - - extensions.resize(numExtensions); - VK_CHECK_CALL_RET(vkEnumerateDeviceExtensionProperties( - device, nullptr, &numExtensions, extensions.data())); - - return VK_SUCCESS; -} - -// Set up the Vulkan device from the SYCL one -VkResult setupDevice(const sycl::device &dev) { - uint32_t physicalDeviceCount = 0; - // Get all physical devices. - VK_CHECK_CALL_RET( - vkEnumeratePhysicalDevices(vk_instance, &physicalDeviceCount, nullptr)); - if (physicalDeviceCount == 0) { - // If no physical devices found, return error. - return VK_ERROR_DEVICE_LOST; - } - std::vector physicalDevices(physicalDeviceCount); - VK_CHECK_CALL_RET(vkEnumeratePhysicalDevices( - vk_instance, &physicalDeviceCount, physicalDevices.data())); - - bool foundDevice = false; - - // Define the required device extensions to run the tests. - static constexpr const char *requiredExtensions[] = { - VK_KHR_GET_MEMORY_REQUIREMENTS_2_EXTENSION_NAME, - VK_KHR_EXTERNAL_MEMORY_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_EXT_EXTERNAL_MEMORY_DMA_BUF_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 - // that also supports all our required device extensions - for (int i = 0; i < physicalDeviceCount; i++) { - vk_physical_device = physicalDevices[i]; - - VkPhysicalDeviceIDProperties devIDProps = { - .sType = VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_ID_PROPERTIES}; - - VkPhysicalDeviceProperties2 devProps2 = { - .sType = VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_PROPERTIES_2, - .pNext = &devIDProps}; - - vkGetPhysicalDeviceProperties2(vk_physical_device, &devProps2); - - if (!std::equal(std::begin(UUID), std::end(UUID), - std::begin(devIDProps.deviceUUID), - std::begin(devIDProps.deviceUUID) + VK_UUID_SIZE)) { - continue; - } - - // Check if the device supports the required extensions. - std::vector supportedDeviceExtensions; - getSupportedDeviceExtensions(supportedDeviceExtensions, vk_physical_device); - const bool hasRequiredExtensions = std::all_of( - std::begin(requiredExtensions), std::end(requiredExtensions), - [&](std::string_view requiredExt) { - auto it = std::find_if(std::begin(supportedDeviceExtensions), - std::end(supportedDeviceExtensions), - [&](const VkExtensionProperties &ext) { - return (ext.extensionName == requiredExt); - }); - 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 (std::string_view optionalExt : optionalExtensions) { - auto it = std::find_if(std::begin(supportedDeviceExtensions), - std::end(supportedDeviceExtensions), - [&](const VkExtensionProperties &ext) { - return (ext.extensionName == optionalExt); - }); - if (it != std::end(supportedDeviceExtensions)) { - enabledDeviceExtensions.push_back(optionalExt.data()); - if (optionalExt == VK_KHR_EXTERNAL_SEMAPHORE_EXTENSION_NAME) { - supportsExternalSemaphore = true; - } - if (optionalExt == VK_EXT_EXTERNAL_MEMORY_DMA_BUF_EXTENSION_NAME) { - supportsDmaBuf = true; - } - } - } - - foundDevice = true; - std::cout << "Found suitable Vulkan device: " - << devProps2.properties.deviceName << std::endl; - break; - } - - // If no device was found that supports all our required extensions return an - // error. - if (!foundDevice) { - std::cerr << "Failed to find suitable device!\n"; - return VK_ERROR_DEVICE_LOST; - } - - // Get queue families and assign queue family indices for compute and transfer - // queues. - uint32_t queueFamilyCount = 0; - vkGetPhysicalDeviceQueueFamilyProperties(vk_physical_device, - &queueFamilyCount, nullptr); - std::vector queueFamilies(queueFamilyCount); - vkGetPhysicalDeviceQueueFamilyProperties( - vk_physical_device, &queueFamilyCount, queueFamilies.data()); - uint32_t i = 0; - bool computeQueueFamilyFound = false; - bool transferQueueFamilyFound = false; - for (auto &qf : queueFamilies) { - // Queue families that support `VK_QUEUE_COMPUTE_BIT` or - // `VK_QUEUE_TRANSFER_BIT` capabilities should also implicitly support - // `VK_QUEUE_GRAPHICS_BIT`. - // `VK_QUEUE_GRAPHICS_BIT` support is required for the `depth_format.cpp` - // test. - if (!computeQueueFamilyFound && (qf.queueFlags & VK_QUEUE_COMPUTE_BIT) && - (qf.queueFlags & VK_QUEUE_GRAPHICS_BIT)) { - vk_computeQueueFamilyIndex = i; - computeQueueFamilyFound = true; - } - if (!transferQueueFamilyFound && (qf.queueFlags & VK_QUEUE_TRANSFER_BIT) && - (qf.queueFlags & VK_QUEUE_GRAPHICS_BIT)) { - vk_transferQueueFamilyIndex = i; - transferQueueFamilyFound = true; - } - ++i; - } - - // Populate queue information prior to Vulkan device creation. - float queuePriority = 1.f; - std::vector qcis; - if (vk_computeQueueFamilyIndex == vk_transferQueueFamilyIndex) { - qcis.resize(1); - qcis[0].sType = VK_STRUCTURE_TYPE_DEVICE_QUEUE_CREATE_INFO; - qcis[0].queueFamilyIndex = vk_transferQueueFamilyIndex; - qcis[0].queueCount = 1; - qcis[0].pQueuePriorities = &queuePriority; - } else { - qcis.resize(2); - qcis[0].sType = VK_STRUCTURE_TYPE_DEVICE_QUEUE_CREATE_INFO; - qcis[0].queueFamilyIndex = vk_transferQueueFamilyIndex; - qcis[0].queueCount = 1; - qcis[0].pQueuePriorities = &queuePriority; - - qcis[1].sType = VK_STRUCTURE_TYPE_DEVICE_QUEUE_CREATE_INFO; - qcis[1].queueFamilyIndex = vk_computeQueueFamilyIndex; - qcis[1].queueCount = 1; - qcis[1].pQueuePriorities = &queuePriority; - } - - VkPhysicalDeviceFeatures deviceFeatures = {}; - - // Create the Vulkan device with the above queues, extensions, and layers. - VkDeviceCreateInfo dci = {}; - dci.sType = VK_STRUCTURE_TYPE_DEVICE_CREATE_INFO; - dci.pQueueCreateInfos = qcis.data(); - dci.queueCreateInfoCount = qcis.size(); - dci.pEnabledFeatures = &deviceFeatures; - dci.enabledExtensionCount = enabledDeviceExtensions.size(); - dci.ppEnabledExtensionNames = enabledDeviceExtensions.data(); - - VK_CHECK_CALL_RET( - vkCreateDevice(vk_physical_device, &dci, nullptr, &vk_device)); - - // Get the Vulkan queues from the device. - vkGetDeviceQueue(vk_device, vk_transferQueueFamilyIndex, 0, - &vk_transfer_queue); - vkGetDeviceQueue(vk_device, vk_computeQueueFamilyIndex, 0, &vk_compute_queue); - - // Get function pointers for memory and semaphore handle exportation. - // Functions will depend on the OS being compiled for. -#ifdef _WIN32 - vk_GetMemoryWin32HandleKHR = - (PFN_vkGetMemoryWin32HandleKHR)vkGetDeviceProcAddr( - vk_device, "vkGetMemoryWin32HandleKHR"); - if (!vk_GetMemoryWin32HandleKHR) { - std::cerr - << "Could not get func pointer to \"vkGetMemoryWin32HandleKHR\"!\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 = - (PFN_vkGetMemoryFdKHR)vkGetDeviceProcAddr(vk_device, "vkGetMemoryFdKHR"); - if (!vk_getMemoryFdKHR) { - std::cerr << "Could not get func pointer to \"vkGetMemoryFdKHR\"!\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 - - vk_getImageMemoryRequirements2 = - reinterpret_cast( - vkGetDeviceProcAddr(vk_device, "vkGetImageMemoryRequirements2KHR")); - if (!vk_getImageMemoryRequirements2) { - std::cerr << "Could not get func pointer to " - "\"vkGetImageMemoryRequirements2KHR\"!\n"; - return VK_ERROR_UNKNOWN; - } - - return VK_SUCCESS; -} - -/* -This function sets up Vulkan command buffers. -Firstly we create command pools for each of the queues that can be used. -We have two queue types which can be used: - - A transfer queue, used for data movement operations - - A compute queue, used for shader invocation operations -We allocate command buffers from these command pools. -Note that some Vulkan instances may provide queues with transfer and compute -capabilities. If this is the case, we only create one command pool, and one -command buffer. -*/ -VkResult setupCommandBuffers() { - VkCommandPoolCreateInfo cpci = {}; - cpci.sType = VK_STRUCTURE_TYPE_COMMAND_POOL_CREATE_INFO; - cpci.queueFamilyIndex = vk_computeQueueFamilyIndex; - cpci.flags = VK_COMMAND_POOL_CREATE_RESET_COMMAND_BUFFER_BIT; - VK_CHECK_CALL_RET( - vkCreateCommandPool(vk_device, &cpci, nullptr, &vk_computeCmdPool)); - - if (vk_computeQueueFamilyIndex == vk_transferQueueFamilyIndex) { - vk_transferCmdPool = vk_computeCmdPool; - } else { - VkCommandPoolCreateInfo cpci = {}; - cpci.sType = VK_STRUCTURE_TYPE_COMMAND_POOL_CREATE_INFO; - cpci.queueFamilyIndex = vk_transferQueueFamilyIndex; - cpci.flags = VK_COMMAND_POOL_CREATE_RESET_COMMAND_BUFFER_BIT; - VK_CHECK_CALL_RET( - vkCreateCommandPool(vk_device, &cpci, nullptr, &vk_transferCmdPool)); - } - - { - VkCommandBufferAllocateInfo cbai = {}; - cbai.sType = VK_STRUCTURE_TYPE_COMMAND_BUFFER_ALLOCATE_INFO; - cbai.commandPool = vk_computeCmdPool; - cbai.level = VK_COMMAND_BUFFER_LEVEL_PRIMARY; - cbai.commandBufferCount = 1; - VK_CHECK_CALL_RET( - vkAllocateCommandBuffers(vk_device, &cbai, &vk_computeCmdBuffer)); - } - - { - VkCommandBufferAllocateInfo cbai = {}; - cbai.sType = VK_STRUCTURE_TYPE_COMMAND_BUFFER_ALLOCATE_INFO; - cbai.commandPool = vk_transferCmdPool; - cbai.level = VK_COMMAND_BUFFER_LEVEL_PRIMARY; - cbai.commandBufferCount = 2; - VK_CHECK_CALL_RET( - vkAllocateCommandBuffers(vk_device, &cbai, vk_transferCmdBuffers)); - } - - return VK_SUCCESS; -} - -/* -Create a Vulkan buffer with a specified size and 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"; - return VK_NULL_HANDLE; - } - return buffer; -} - -/* -Create a Vulkan image with a specified image type, format, extent, and usage. -This function also allows users to specify whether the image will be exportable, -in which case the appropriate extension struct is populated based on the OS the -program is compiled for. -*/ -VkImage createImage(VkImageType type, VkFormat format, VkExtent3D extent, - VkImageUsageFlags usage, size_t mipLevels, - bool linearTiling = false, bool exportable = true) { - VkImageCreateInfo ici = {}; - ici.sType = VK_STRUCTURE_TYPE_IMAGE_CREATE_INFO; - ici.imageType = type; - ici.format = format; - ici.extent = extent; - ici.mipLevels = mipLevels; - ici.arrayLayers = 1; - ici.usage = usage; - ici.sharingMode = VK_SHARING_MODE_EXCLUSIVE; - ici.samples = VK_SAMPLE_COUNT_1_BIT; - - if (linearTiling) { - ici.tiling = VK_IMAGE_TILING_LINEAR; - } - - VkExternalMemoryImageCreateInfo emici = {}; - if (exportable) { - emici.sType = VK_STRUCTURE_TYPE_EXTERNAL_MEMORY_IMAGE_CREATE_INFO; -#ifdef _WIN32 - emici.handleTypes = VK_EXTERNAL_MEMORY_HANDLE_TYPE_OPAQUE_WIN32_BIT; -#else - emici.handleTypes = VK_EXTERNAL_MEMORY_HANDLE_TYPE_OPAQUE_FD_BIT; -#endif - - ici.pNext = &emici; - } - - VkImage image; - if (vkCreateImage(vk_device, &ici, nullptr, &image)) { - std::cerr << "Could not create image!\n"; - return VK_NULL_HANDLE; - } - return image; -} - -/* -Returns the row pitch with which a linear image's first subresource was -created with in bytes. -*/ -VkDeviceSize getImageRowPitch(VkImage image) { - - VkImageSubresource imageSubresource = {}; - imageSubresource.aspectMask = VK_IMAGE_ASPECT_COLOR_BIT; - imageSubresource.mipLevel = 0; - imageSubresource.arrayLayer = 0; - - VkSubresourceLayout subresourceLayout = {}; - vkGetImageSubresourceLayout(vk_device, image, &imageSubresource, - &subresourceLayout); - - return subresourceLayout.rowPitch; -} - -/* -Allocate `size` of device memory of the specified memory type. -This function also allows users to specify whether the memory will be -exportable, in which case the appropriate extension struct is populated based on -the OS the program is compiled for. -*/ -VkDeviceMemory allocateDeviceMemory(size_t size, uint32_t memoryTypeIndex, - VkImage image, bool exportable = true) { - VkMemoryAllocateInfo mai{}; - mai.sType = VK_STRUCTURE_TYPE_MEMORY_ALLOCATE_INFO; - mai.allocationSize = size; - mai.memoryTypeIndex = memoryTypeIndex; - - // A tiled image's real backing size is padded beyond width*height*texel and - // is only known from vkGetImageMemoryRequirements; the caller's element-count - // based `size` under-allocates on drivers that pad (e.g. DG2). Grow the - // allocation to the image requirement so the bind is spec-valid and the - // exported handle describes the whole image. - if (image != VK_NULL_HANDLE) { - VkMemoryRequirements imageReq{}; - vkGetImageMemoryRequirements(vk_device, image, &imageReq); - if (imageReq.size > mai.allocationSize) - mai.allocationSize = imageReq.size; - } - - // Use dedicated allocation whenever backing an exportable image, not only - // when the driver reports it as required (some drivers under-report this for - // external memory). - const bool useDedicated = - requiresDedicatedAllocation || (exportable && image != VK_NULL_HANDLE); - - VkMemoryDedicatedAllocateInfoKHR dedicatedInfo{}; - if (useDedicated) { - dedicatedInfo.sType = VK_STRUCTURE_TYPE_MEMORY_DEDICATED_ALLOCATE_INFO_KHR; - dedicatedInfo.image = image; - dedicatedInfo.buffer = VK_NULL_HANDLE; - mai.pNext = &dedicatedInfo; - } - - VkExportMemoryAllocateInfo emai{}; - if (exportable) { - emai.sType = VK_STRUCTURE_TYPE_EXPORT_MEMORY_ALLOCATE_INFO; -#ifdef _WIN32 - emai.handleTypes = VK_EXTERNAL_MEMORY_HANDLE_TYPE_OPAQUE_WIN32_BIT; -#else - emai.handleTypes = VK_EXTERNAL_MEMORY_HANDLE_TYPE_OPAQUE_FD_BIT; -#endif - if (useDedicated) { - dedicatedInfo.pNext = &emai; - } else { - mai.pNext = &emai; - } - } - - VkDeviceMemory memory; - if (vkAllocateMemory(vk_device, &mai, nullptr, &memory) != VK_SUCCESS) { - std::cerr << "Could not allocate device memory!\n"; - return VK_NULL_HANDLE; - } - - return memory; -} - -/* -Import device memory from a file descriptor. -*/ -template -VkDeviceMemory importDeviceMemory(size_t size, uint32_t memoryTypeIndex, - InteropHandleT fd); - -#ifndef _WIN32 -// Specialization for importing device memory from file descriptors. -template <> -VkDeviceMemory importDeviceMemory(size_t size, uint32_t memoryTypeIndex, - int fd) { - - VkMemoryAllocateInfo mai{}; - mai.sType = VK_STRUCTURE_TYPE_MEMORY_ALLOCATE_INFO; - mai.allocationSize = size; - mai.memoryTypeIndex = memoryTypeIndex; - - VkImportMemoryFdInfoKHR importInfo{}; - importInfo.sType = VK_STRUCTURE_TYPE_IMPORT_MEMORY_FD_INFO_KHR; - importInfo.handleType = VK_EXTERNAL_MEMORY_HANDLE_TYPE_OPAQUE_FD_BIT; - importInfo.fd = fd; - mai.pNext = &importInfo; - - VkDeviceMemory memory; - if (vkAllocateMemory(vk_device, &mai, nullptr, &memory) != VK_SUCCESS) { - std::cerr << "importDeviceMemoryFD -- Could not allocate device memory!\n"; - return VK_NULL_HANDLE; - } - - return memory; -} - -#else - -// Specialization for importing device memory from win32 handles. -template <> -VkDeviceMemory importDeviceMemory(size_t size, uint32_t memoryTypeIndex, - void *winHandle) { - - VkMemoryAllocateInfo mai{}; - mai.sType = VK_STRUCTURE_TYPE_MEMORY_ALLOCATE_INFO; - mai.allocationSize = size; - mai.memoryTypeIndex = memoryTypeIndex; - - VkImportMemoryWin32HandleInfoKHR importInfo{}; - importInfo.sType = VK_STRUCTURE_TYPE_IMPORT_MEMORY_WIN32_HANDLE_INFO_KHR; - importInfo.handleType = VK_EXTERNAL_MEMORY_HANDLE_TYPE_OPAQUE_WIN32_BIT; - importInfo.handle = winHandle; - mai.pNext = &importInfo; - - VkDeviceMemory memory; - if (vkAllocateMemory(vk_device, &mai, nullptr, &memory) != VK_SUCCESS) { - std::cerr << "importDeviceMemoryFD -- Could not allocate device memory!\n"; - return VK_NULL_HANDLE; - } - - return memory; -} - -#endif // _WIN32 - -/* -Retrieve the image memory type index for the Vulkan device based on the memory -property flags passed. -*/ -uint32_t getImageMemoryTypeIndex(VkImage image, VkMemoryPropertyFlags flags, - VkMemoryRequirements &memRequirements) { - VkMemoryRequirements2 memoryRequirements2{}; - memoryRequirements2.sType = VK_STRUCTURE_TYPE_MEMORY_REQUIREMENTS_2; - - VkMemoryDedicatedRequirements dedicatedRequirements{}; - if (supportsDedicatedAllocation) { - dedicatedRequirements.sType = - VK_STRUCTURE_TYPE_MEMORY_DEDICATED_REQUIREMENTS; - memoryRequirements2.pNext = &dedicatedRequirements; - } - - VkImageMemoryRequirementsInfo2 imageRequirementsInfo{}; - imageRequirementsInfo.sType = - VK_STRUCTURE_TYPE_IMAGE_MEMORY_REQUIREMENTS_INFO_2; - imageRequirementsInfo.image = image; - - vk_getImageMemoryRequirements2(vk_device, &imageRequirementsInfo, - &memoryRequirements2); - - if (dedicatedRequirements.requiresDedicatedAllocation) { - requiresDedicatedAllocation = true; - } - - VkPhysicalDeviceMemoryProperties memProperties; - vkGetPhysicalDeviceMemoryProperties(vk_physical_device, &memProperties); - - memRequirements = memoryRequirements2.memoryRequirements; - for (uint32_t i = 0; i < memProperties.memoryTypeCount; i++) { - if ((memRequirements.memoryTypeBits & (1 << i)) && - (memProperties.memoryTypes[i].propertyFlags & flags) == flags) { - return i; - } - } - std::cerr << "Image memory type index not found!\n"; - return 0; -} - -/* -Retrieve the buffer memory type index for the Vulkan device based on the memory -property flags passed. -*/ -uint32_t getBufferMemoryTypeIndex(VkBuffer buffer, - VkMemoryPropertyFlags flags) { - VkMemoryRequirements memRequirements; - vkGetBufferMemoryRequirements(vk_device, buffer, &memRequirements); - - VkPhysicalDeviceMemoryProperties memProperties; - vkGetPhysicalDeviceMemoryProperties(vk_physical_device, &memProperties); - - for (uint32_t i = 0; i < memProperties.memoryTypeCount; i++) { - if ((memRequirements.memoryTypeBits & (1 << i)) && - (memProperties.memoryTypes[i].propertyFlags & flags) == flags) { - return i; - } - } - std::cerr << "Buffer memory type index not found!\n"; - return 0; -} - -/* -Destroy Vulkan objects. -This function is called towards the end of Vulkan program execution. -*/ -VkResult cleanup() { - - if (vk_computeQueueFamilyIndex == vk_transferQueueFamilyIndex) { - vkDestroyCommandPool(vk_device, vk_computeCmdPool, nullptr); - } else { - vkDestroyCommandPool(vk_device, vk_computeCmdPool, nullptr); - vkDestroyCommandPool(vk_device, vk_transferCmdPool, nullptr); - } - - auto destroyDebugUtilsMessenger = - (PFN_vkDestroyDebugUtilsMessengerEXT)vkGetInstanceProcAddr( - vk_instance, "vkDestroyDebugUtilsMessengerEXT"); - if (destroyDebugUtilsMessenger != nullptr) { - destroyDebugUtilsMessenger(vk_instance, vk_debug_messenger, nullptr); - } - vkDestroyDevice(vk_device, nullptr); - vkDestroyInstance(vk_instance, nullptr); - return VK_SUCCESS; -} - -#ifdef _WIN32 - -/* -Retrieve a win32 memory handle for a given Vulkan device memory allocation. -*/ -HANDLE getMemoryWin32Handle(VkDeviceMemory memory) { - - HANDLE retHandle = 0; - - VkMemoryGetWin32HandleInfoKHR mgwhi = {}; - mgwhi.sType = VK_STRUCTURE_TYPE_MEMORY_GET_WIN32_HANDLE_INFO_KHR; - mgwhi.memory = memory; - mgwhi.handleType = VK_EXTERNAL_MEMORY_HANDLE_TYPE_OPAQUE_WIN32_BIT; - - if (vk_GetMemoryWin32HandleKHR != nullptr) { - VK_CHECK_CALL(vk_GetMemoryWin32HandleKHR(vk_device, &mgwhi, &retHandle)); - } else { - std::cerr << "Could not get win32 handle!\n"; - return 0; - } - - return retHandle; -} - -/* -Retrieve a win32 memory handle for a given Vulkan semaphore object. -*/ -HANDLE getSemaphoreWin32Handle(VkSemaphore semaphore) { - - HANDLE retHandle = 0; - - VkSemaphoreGetWin32HandleInfoKHR sghwi = {}; - sghwi.sType = VK_STRUCTURE_TYPE_SEMAPHORE_GET_WIN32_HANDLE_INFO_KHR; - 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 { - std::cerr << "Could not get semaphore opaque file descriptor!\n"; - return 0; - } - - return retHandle; -} - -#else - -/* -Retrieve an opaque file descriptor handle for a given Vulkan memory allocation. -*/ -template -int getMemoryOpaqueFD(VkDeviceMemory memory) { - VkMemoryGetFdInfoKHR mgfi = {}; - mgfi.sType = VK_STRUCTURE_TYPE_MEMORY_GET_FD_INFO_KHR; - mgfi.memory = memory; - mgfi.handleType = vulkanHandleType; - - int fd = 0; - if (vk_getMemoryFdKHR != nullptr) { - VK_CHECK_CALL(vk_getMemoryFdKHR(vk_device, &mgfi, &fd)); - } else { - std::cerr << "Could not get memory opaque file descriptor!\n"; - return 0; - } - - return fd; -} - -/* -Retrieve an opaque file descriptor handle for a given Vulkan semaphore object. -*/ -int getSemaphoreOpaqueFD(VkSemaphore semaphore) { - VkSemaphoreGetFdInfoKHR sgfi = {}; - sgfi.sType = VK_STRUCTURE_TYPE_SEMAPHORE_GET_FD_INFO_KHR; - sgfi.semaphore = 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 { - std::cerr << "Could not get semaphore opaque file descriptor!\n"; - return 0; - } - - return fd; -} -#endif - -/* -Populate a generic image memory barrier for a specific Vulkan image. -This function assumes we are transitioning from an undefined image layout to a -general image layout, which is sufficient for our current Vulkan tests. -*/ -auto createImageMemoryBarrier(VkImage &img, size_t mipLevels) { - VkImageMemoryBarrier barrierInput = {}; - barrierInput.sType = VK_STRUCTURE_TYPE_IMAGE_MEMORY_BARRIER; - barrierInput.oldLayout = VK_IMAGE_LAYOUT_UNDEFINED; - barrierInput.newLayout = VK_IMAGE_LAYOUT_GENERAL; - barrierInput.srcQueueFamilyIndex = VK_QUEUE_FAMILY_IGNORED; - barrierInput.dstQueueFamilyIndex = VK_QUEUE_FAMILY_IGNORED; - barrierInput.image = img; - barrierInput.subresourceRange.aspectMask = VK_IMAGE_ASPECT_COLOR_BIT; - barrierInput.subresourceRange.levelCount = mipLevels; - barrierInput.subresourceRange.layerCount = 1; - barrierInput.srcAccessMask = 0; - barrierInput.dstAccessMask = VK_ACCESS_TRANSFER_WRITE_BIT; - return barrierInput; -} - -/* -This struct contains Vulkan resources used in test files, and is used to -simplify the code within these tests. -The constructor creates images, allocates device memory required for those -images, and binds that memory to the created image. -The destructor cleans up the memory allocations and destroys the image and -staging buffer used to transfer data to that image. -*/ -struct vulkan_image_test_resources_t { - VkImage vkImage; - VkDeviceMemory imageMemory; - VkBuffer stagingBuffer; - VkDeviceMemory stagingMemory; - - 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, 1); - VkMemoryRequirements memRequirements; - auto inputImageMemoryTypeIndex = vkutil::getImageMemoryTypeIndex( - vkImage, VK_MEMORY_PROPERTY_DEVICE_LOCAL_BIT, memRequirements); - imageMemory = vkutil::allocateDeviceMemory( - imageSizeBytes, inputImageMemoryTypeIndex, vkImage); - VK_CHECK_CALL( - vkBindImageMemory(vk_device, vkImage, imageMemory, 0 /*memoryOffset*/)); - - stagingBuffer = vkutil::createBuffer(imageSizeBytes, - VK_BUFFER_USAGE_TRANSFER_SRC_BIT | - VK_BUFFER_USAGE_TRANSFER_DST_BIT); - auto inputStagingMemoryTypeIndex = vkutil::getBufferMemoryTypeIndex( - stagingBuffer, VK_MEMORY_PROPERTY_HOST_VISIBLE_BIT | - VK_MEMORY_PROPERTY_HOST_COHERENT_BIT); - stagingMemory = vkutil::allocateDeviceMemory( - imageSizeBytes, inputStagingMemoryTypeIndex, nullptr /*image*/, - false /*exportable*/); - VK_CHECK_CALL(vkBindBufferMemory(vk_device, stagingBuffer, stagingMemory, - 0 /*memoryOffset*/)); - } - - ~vulkan_image_test_resources_t() { - vkDestroyBuffer(vk_device, stagingBuffer, nullptr); - vkDestroyImage(vk_device, vkImage, nullptr); - vkFreeMemory(vk_device, stagingMemory, nullptr); - vkFreeMemory(vk_device, imageMemory, nullptr); - } -}; - -/* -Convert a SYCL image channel order and image channel type to a corresponding -Vulkan format. -*/ -VkFormat to_vulkan_format(sycl::image_channel_order order, - sycl::image_channel_type channel_type) { - if (channel_type == sycl::image_channel_type::unorm_int8) { - switch (order) { - case sycl::image_channel_order::r: - return VK_FORMAT_R8_UNORM; - case sycl::image_channel_order::rg: - return VK_FORMAT_R8G8_UNORM; - case sycl::image_channel_order::rgb: - return VK_FORMAT_R8G8B8_UNORM; - case sycl::image_channel_order::rgba: - return VK_FORMAT_R8G8B8A8_UNORM; - default: { - std::cerr << "error in converting to vulkan format\n"; - exit(-1); - } - } - } else if (channel_type == sycl::image_channel_type::signed_int8) { - - switch (order) { - case sycl::image_channel_order::r: - return VK_FORMAT_R8_SINT; - case sycl::image_channel_order::rg: - return VK_FORMAT_R8G8_SINT; - case sycl::image_channel_order::rgba: - return VK_FORMAT_R8G8B8A8_SINT; - default: { - std::cerr << "error in converting to vulkan format\n"; - exit(-1); - } - } - } else if (channel_type == sycl::image_channel_type::unsigned_int32) { - - switch (order) { - case sycl::image_channel_order::r: - return VK_FORMAT_R32_UINT; - case sycl::image_channel_order::rg: - return VK_FORMAT_R32G32_UINT; - case sycl::image_channel_order::rgba: - return VK_FORMAT_R32G32B32A32_UINT; - default: { - std::cerr << "error in converting to vulkan format\n"; - exit(-1); - } - } - } else if (channel_type == sycl::image_channel_type::signed_int32) { - switch (order) { - case sycl::image_channel_order::r: - return VK_FORMAT_R32_SINT; - case sycl::image_channel_order::rg: - return VK_FORMAT_R32G32_SINT; - case sycl::image_channel_order::rgba: - return VK_FORMAT_R32G32B32A32_SINT; - default: { - std::cerr << "error in converting to vulkan format\n"; - exit(-1); - } - } - } else if (channel_type == sycl::image_channel_type::signed_int16) { - switch (order) { - case sycl::image_channel_order::r: - return VK_FORMAT_R16_SINT; - case sycl::image_channel_order::rg: - return VK_FORMAT_R16G16_SINT; - case sycl::image_channel_order::rgba: - return VK_FORMAT_R16G16B16A16_SINT; - default: { - std::cerr << "error in converting to vulkan format\n"; - exit(-1); - } - } - } else if (channel_type == sycl::image_channel_type::fp16) { - switch (order) { - case sycl::image_channel_order::r: - return VK_FORMAT_R16_SFLOAT; - case sycl::image_channel_order::rg: - return VK_FORMAT_R16G16_SFLOAT; - case sycl::image_channel_order::rgb: - return VK_FORMAT_R16G16B16_SFLOAT; - case sycl::image_channel_order::rgba: - return VK_FORMAT_R16G16B16A16_SFLOAT; - default: { - std::cerr << "error in converting to vulkan format\n"; - exit(-1); - } - } - } else if (channel_type == sycl::image_channel_type::fp32) { - switch (order) { - case sycl::image_channel_order::r: - return VK_FORMAT_R32_SFLOAT; - case sycl::image_channel_order::rg: - return VK_FORMAT_R32G32_SFLOAT; - case sycl::image_channel_order::rgba: - return VK_FORMAT_R32G32B32A32_SFLOAT; - default: { - std::cerr << "error in converting to vulkan format\n"; - exit(-1); - } - } - } else { - std::cerr - << "error in converting to vulkan format - channel type not included\n"; - exit(-1); - } -} - -} // namespace vkutil - -namespace util { - -template -bool is_equal(DType lhs, DType rhs, float epsilon = 0.0001f) { - if constexpr (std::is_floating_point_v) { - return (std::abs(lhs - rhs) < epsilon); - } else { - return lhs == rhs; - } -} - -} // namespace util diff --git a/sycl/test-e2e/MemoryExport/export_memory_to_vulkan.cpp b/sycl/test-e2e/MemoryExport/export_memory_to_vulkan.cpp index 3ff2fb3e1ab0d..f7d86a527ec9e 100644 --- a/sycl/test-e2e/MemoryExport/export_memory_to_vulkan.cpp +++ b/sycl/test-e2e/MemoryExport/export_memory_to_vulkan.cpp @@ -20,7 +20,7 @@ #include #include -#include "../CommonUtils/vulkan_common.hpp" +#include "../bindless_images/vulkan_interop/sycl_vulkan_setup.hpp" namespace syclexp = sycl::ext::oneapi::experimental; @@ -80,7 +80,8 @@ void cleanupSycl(const sycl::device &SyclDevice) { SyclContext); } -int runTest(sycl::device &SyclDevice, const size_t MemorySizeBytes) { +int runTest(VulkanContext &VulkanCtx, sycl::device &SyclDevice, + const size_t MemorySizeBytes) { sycl::context SyclContext = sycl::context(SyclDevice); sycl::queue SyclQueue(SyclContext, SyclDevice); @@ -89,39 +90,52 @@ int runTest(sycl::device &SyclDevice, const size_t MemorySizeBytes) { VkDeviceMemory VkImportedBufferMemory; { - VkImportedBuffer = vkutil::createBuffer( - MemorySizeBytes, - VK_BUFFER_USAGE_TRANSFER_SRC_BIT | VK_BUFFER_USAGE_TRANSFER_DST_BIT | - VK_BUFFER_USAGE_STORAGE_BUFFER_BIT, - true /*exportable*/); - auto InputBufferMemTypeIndex = vkutil::getBufferMemoryTypeIndex( - VkImportedBuffer, VK_MEMORY_PROPERTY_DEVICE_LOCAL_BIT); - - VkImportedBufferMemory = vkutil::importDeviceMemory( - MemorySizeBytes, InputBufferMemTypeIndex, ExportableMemoryHandle); - - VK_CHECK_CALL(vkBindBufferMemory(vk_device, VkImportedBuffer, - VkImportedBufferMemory, - 0 /*memoryOffset*/)); + VkExternalMemoryBufferCreateInfo ExternalInfo = { + VK_STRUCTURE_TYPE_EXTERNAL_MEMORY_BUFFER_CREATE_INFO}; + ExternalInfo.handleTypes = PLATFORM_MEM_HANDLE_TYPE; + VkBufferCreateInfo BufferInfo = {VK_STRUCTURE_TYPE_BUFFER_CREATE_INFO}; + BufferInfo.pNext = &ExternalInfo; + BufferInfo.size = MemorySizeBytes; + BufferInfo.usage = VK_BUFFER_USAGE_TRANSFER_SRC_BIT | + VK_BUFFER_USAGE_TRANSFER_DST_BIT | + VK_BUFFER_USAGE_STORAGE_BUFFER_BIT; + BufferInfo.sharingMode = VK_SHARING_MODE_EXCLUSIVE; + VK_CHECK(vkCreateBuffer(VulkanCtx.device, &BufferInfo, nullptr, + &VkImportedBuffer)); + + VkMemoryRequirements Requirements; + vkGetBufferMemoryRequirements(VulkanCtx.device, VkImportedBuffer, + &Requirements); + VkMemoryAllocateInfo AllocateInfo = { + VK_STRUCTURE_TYPE_MEMORY_ALLOCATE_INFO}; + AllocateInfo.allocationSize = Requirements.size; + AllocateInfo.memoryTypeIndex = + findMemoryType(VulkanCtx.physicalDevice, Requirements.memoryTypeBits, + VK_MEMORY_PROPERTY_DEVICE_LOCAL_BIT); +#ifdef _WIN32 + VkImportMemoryWin32HandleInfoKHR ImportInfo = { + VK_STRUCTURE_TYPE_IMPORT_MEMORY_WIN32_HANDLE_INFO_KHR}; + ImportInfo.handleType = PLATFORM_MEM_HANDLE_TYPE; + ImportInfo.handle = ExportableMemoryHandle; +#else + VkImportMemoryFdInfoKHR ImportInfo = { + VK_STRUCTURE_TYPE_IMPORT_MEMORY_FD_INFO_KHR}; + ImportInfo.handleType = PLATFORM_MEM_HANDLE_TYPE; + ImportInfo.fd = ExportableMemoryHandle; +#endif + AllocateInfo.pNext = &ImportInfo; + VK_CHECK(vkAllocateMemory(VulkanCtx.device, &AllocateInfo, nullptr, + &VkImportedBufferMemory)); + VK_CHECK(vkBindBufferMemory(VulkanCtx.device, VkImportedBuffer, + VkImportedBufferMemory, 0)); } // Allocate temporary staging buffer and copy imported data to host. VulkanOutput.resize(MemorySizeBytes / sizeof(DataT), 0); { - VkBuffer StagingBuffer; - VkDeviceMemory StagingMemory; - - StagingBuffer = vkutil::createBuffer(MemorySizeBytes, - 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( - MemorySizeBytes, InputStagingMemTypeIndex, VK_NULL_HANDLE /*image*/, - false /*exportable*/); - VK_CHECK_CALL(vkBindBufferMemory(vk_device, StagingBuffer, StagingMemory, - 0 /*memoryOffset*/)); + auto Staging = createStagingBuffer(VulkanCtx, MemorySizeBytes, + VK_BUFFER_USAGE_TRANSFER_SRC_BIT | + VK_BUFFER_USAGE_TRANSFER_DST_BIT); // Copy imported buffer to host visible staging buffer. VkCommandBufferBeginInfo Cbbi = {}; @@ -131,40 +145,28 @@ int runTest(sycl::device &SyclDevice, const size_t MemorySizeBytes) { VkBufferCopy CopyRegion = {}; CopyRegion.size = MemorySizeBytes; - VK_CHECK_CALL(vkBeginCommandBuffer(vk_transferCmdBuffers[0], &Cbbi)); - vkCmdCopyBuffer(vk_transferCmdBuffers[0], VkImportedBuffer, StagingBuffer, + VkCommandPool Pool; + VkCommandBuffer CommandBuffer = createCommandBuffer(VulkanCtx, Pool); + VK_CHECK(vkBeginCommandBuffer(CommandBuffer, &Cbbi)); + vkCmdCopyBuffer(CommandBuffer, VkImportedBuffer, Staging.buffer, 1 /*regionCount*/, &CopyRegion); - 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)); + submitCommandBuffer(VulkanCtx, CommandBuffer, Pool); // Copy host visible staging buffer data to host. DataT *StagingData = nullptr; - VK_CHECK_CALL(vkMapMemory(vk_device, StagingMemory, 0 /*offset*/, - MemorySizeBytes, 0 /*flags*/, - (void **)&StagingData)); + VK_CHECK(vkMapMemory(VulkanCtx.device, Staging.memory, 0 /*offset*/, + MemorySizeBytes, 0 /*flags*/, (void **)&StagingData)); for (int i = 0; i < MemorySizeBytes / sizeof(DataT); ++i) { VulkanOutput[i] = StagingData[i]; } - vkUnmapMemory(vk_device, StagingMemory); + vkUnmapMemory(VulkanCtx.device, Staging.memory); // Destroy temporary staging buffer and free memory. - vkDestroyBuffer(vk_device, StagingBuffer, nullptr); - vkFreeMemory(vk_device, StagingMemory, nullptr); + cleanupBuffer(VulkanCtx, Staging); } - vkDestroyBuffer(vk_device, VkImportedBuffer, nullptr); - vkFreeMemory(vk_device, VkImportedBufferMemory, nullptr); + vkDestroyBuffer(VulkanCtx.device, VkImportedBuffer, nullptr); + vkFreeMemory(VulkanCtx.device, VkImportedBufferMemory, nullptr); // Print the SYCL imported data. bool Validated = true; @@ -221,45 +223,63 @@ int main(int argc, char *argv[]) { return 3; } - // Init Vulkan. - if (vkutil::setupInstance() != VK_SUCCESS) { - std::cerr << "Instance setup failed!\n"; - return 4; - } - - if (vkutil::setupDevice(SyclDevice) != VK_SUCCESS) { - std::cerr << "Device setup failed!\n"; - return 5; - } - - if (vkutil::setupCommandBuffers() != VK_SUCCESS) { - std::cerr << "Command buffers setup failed!\n"; - return 6; - } - - auto TestPassed = runTest(SyclDevice, MemorySizeBytes); - - if (vkutil::cleanup() != VK_SUCCESS) { - std::cerr << "Cleanup failed!\n"; - return 7; - } + // CleanupFailed is set by the guard's destructor, which only logs on + // failure; capturing the whole test body in a lambda lets us observe + // that flag (after the guard has already run) before main returns. + bool CleanupFailed = false; + + int TestExitCode = [&SyclDevice, &CleanupFailed, MemorySizeBytes]() -> int { + struct SyclCleanupGuard { + const sycl::device &device; + bool &Failed; + ~SyclCleanupGuard() { + try { + cleanupSycl(device); + } catch (const sycl::exception &e) { + std::cerr << "SYCL cleanup failed: " << e.what() << "\n"; + Failed = true; + } catch (...) { + std::cerr << "Unknown exception during SYCL cleanup.\n"; + Failed = true; + } + } + } syclCleanupGuard{SyclDevice, CleanupFailed}; + + // Init Vulkan. + VulkanContext VulkanCtx; + try { + VulkanCtx = createSyclVulkanContext(SyclDevice); + struct VulkanContextGuard { + VulkanContext &context; + ~VulkanContextGuard() { cleanupVulkanContext(context); } + } vulkanContextGuard{VulkanCtx}; + + try { + auto TestPassed = runTest(VulkanCtx, SyclDevice, MemorySizeBytes); + if (TestPassed) { + std::cout << "Test passed!\n"; + return 0; + } + } catch (const std::exception &e) { + std::cerr << "Vulkan test failed: " << e.what() << "\n"; + return 11; + } catch (...) { + std::cerr << "Unknown exception during Vulkan test.\n"; + return 12; + } + } catch (const std::exception &e) { + std::cerr << "Vulkan setup failed: " << e.what() << "\n"; + return 4; + } - // Cleanup SYCL. - try { - cleanupSycl(SyclDevice); - } catch (const sycl::exception &e) { - std::cerr << "SYCL exception caught: " << e.what() << "\n"; - return 8; - } catch (...) { - std::cerr << "Unknown exception caught.\n"; - return 9; - } + std::cerr << "Test failed\n"; + return 10; + }(); - if (TestPassed) { - std::cout << "Test passed!\n"; - return 0; + if (CleanupFailed && TestExitCode == 0) { + std::cerr << "Test failed due to SYCL cleanup error\n"; + return 13; } - std::cerr << "Test failed\n"; - return 10; + return TestExitCode; } 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 2fec4b463a648..4b4a6b62dc3d2 100644 --- a/sycl/test-e2e/bindless_images/vulkan_interop/depth_format.cpp +++ b/sycl/test-e2e/bindless_images/vulkan_interop/depth_format.cpp @@ -12,8 +12,8 @@ // #define VERBOSE_PRINT #include -#include "../../CommonUtils/vulkan_common.hpp" #include "../helpers/common.hpp" +#include "sycl_vulkan_setup.hpp" #include #include @@ -104,8 +104,8 @@ void runSycl(const sycl::device &syclDevice, sycl::range<2> globalSize, } } -bool runTest(const sycl::device &syclDevice, sycl::range<2> dims, - sycl::range<2> localSize) { +bool runTest(VulkanContext &vkCtx, const sycl::device &syclDevice, + sycl::range<2> dims, sycl::range<2> localSize) { const uint32_t imgWidth = static_cast(dims[0]); const uint32_t imgHeight = static_cast(dims[1]); @@ -118,10 +118,8 @@ bool runTest(const sycl::device &syclDevice, sycl::range<2> dims, const VkExtent3D imgExtent = {imgWidth, imgHeight, 1}; - VkImage vkInputImage; - VkDeviceMemory vkInputImageMemory; - VkImage vkOutputImage; - VkDeviceMemory vkOutputImageMemory; + ImageResources inputImage; + ImageResources outputImage; // Real import size; set to the image memory requirement below. size_t importSizeBytes = imgSizeBytes; @@ -137,45 +135,31 @@ bool runTest(const sycl::device &syclDevice, sycl::range<2> dims, { // STORAGE_BIT: SYCL reads/writes this as a storage image; without it the // layout is transfer-only and imported reads land at the wrong offset. - vkInputImage = vkutil::createImage(imgType, imgInFormat, imgExtent, - VK_IMAGE_USAGE_STORAGE_BIT | - VK_IMAGE_USAGE_TRANSFER_SRC_BIT | - VK_IMAGE_USAGE_TRANSFER_DST_BIT, - 1 /*mipLevels*/); + inputImage = createExportableImage( + vkCtx, imgExtent, imgInFormat, imgType, VK_IMAGE_TILING_OPTIMAL, + VK_IMAGE_USAGE_STORAGE_BIT | VK_IMAGE_USAGE_TRANSFER_SRC_BIT | + VK_IMAGE_USAGE_TRANSFER_DST_BIT); VkMemoryRequirements memRequirements; - auto inputImageMemoryTypeIndex = vkutil::getImageMemoryTypeIndex( - vkInputImage, VK_MEMORY_PROPERTY_DEVICE_LOCAL_BIT, memRequirements); + vkGetImageMemoryRequirements(vkCtx.device, inputImage.image, + &memRequirements); // Import must describe the whole (padded) allocation the driver requires. importSizeBytes = std::max(imgSizeBytes, memRequirements.size); - vkInputImageMemory = vkutil::allocateDeviceMemory( - imgSizeBytes, inputImageMemoryTypeIndex, vkInputImage); - VK_CHECK_CALL(vkBindImageMemory(vk_device, vkInputImage, vkInputImageMemory, - 0 /*memoryOffset*/)); // STORAGE_BIT: same as input image; the kernel writes it as a storage // image. - vkOutputImage = vkutil::createImage(imgType, imgOutFormat, imgExtent, - VK_IMAGE_USAGE_STORAGE_BIT | - VK_IMAGE_USAGE_TRANSFER_SRC_BIT | - VK_IMAGE_USAGE_TRANSFER_DST_BIT, - 1 /*mipLevels*/); - VkMemoryRequirements outputMemRequirements; - auto outputImageMemoryTypeIndex = vkutil::getImageMemoryTypeIndex( - vkOutputImage, VK_MEMORY_PROPERTY_DEVICE_LOCAL_BIT, - outputMemRequirements); - vkOutputImageMemory = vkutil::allocateDeviceMemory( - imgSizeBytes, outputImageMemoryTypeIndex, vkOutputImage); - VK_CHECK_CALL(vkBindImageMemory(vk_device, vkOutputImage, - vkOutputImageMemory, 0 /*memoryOffset*/)); + outputImage = createExportableImage( + vkCtx, imgExtent, imgOutFormat, imgType, VK_IMAGE_TILING_OPTIMAL, + VK_IMAGE_USAGE_STORAGE_BIT | VK_IMAGE_USAGE_TRANSFER_SRC_BIT | + VK_IMAGE_USAGE_TRANSFER_DST_BIT); } // Transition image layouts. - printString("Submitting image layout transition\n"); + std::cout << "Submitting image layout transition\n"; { VkImageMemoryBarrier imgInBarrier = - vkutil::createImageMemoryBarrier(vkInputImage, 1 /*mipLevels*/); + createImageMemoryBarrier(inputImage.image, 1); VkImageMemoryBarrier imgOutBarrier = - vkutil::createImageMemoryBarrier(vkOutputImage, 1 /*mipLevels*/); + createImageMemoryBarrier(outputImage.image, 1); // Update aspect mask for the images to VK_IMAGE_ASPECT_DEPTH_BIT. imgInBarrier.subresourceRange.aspectMask = VK_IMAGE_ASPECT_DEPTH_BIT; @@ -185,53 +169,35 @@ bool runTest(const sycl::device &syclDevice, sycl::range<2> dims, cbbi.sType = VK_STRUCTURE_TYPE_COMMAND_BUFFER_BEGIN_INFO; cbbi.flags = VK_COMMAND_BUFFER_USAGE_ONE_TIME_SUBMIT_BIT; - VK_CHECK_CALL(vkBeginCommandBuffer(vk_computeCmdBuffer, &cbbi)); - vkCmdPipelineBarrier(vk_computeCmdBuffer, VK_PIPELINE_STAGE_TOP_OF_PIPE_BIT, + VkCommandPool pool; + VkCommandBuffer commandBuffer = createCommandBuffer(vkCtx, pool); + VK_CHECK(vkBeginCommandBuffer(commandBuffer, &cbbi)); + vkCmdPipelineBarrier(commandBuffer, VK_PIPELINE_STAGE_TOP_OF_PIPE_BIT, VK_PIPELINE_STAGE_TRANSFER_BIT, 0, 0, nullptr, 0, nullptr, 1, &imgInBarrier); - vkCmdPipelineBarrier(vk_computeCmdBuffer, VK_PIPELINE_STAGE_TOP_OF_PIPE_BIT, + vkCmdPipelineBarrier(commandBuffer, VK_PIPELINE_STAGE_TOP_OF_PIPE_BIT, VK_PIPELINE_STAGE_TRANSFER_BIT, 0, 0, nullptr, 0, nullptr, 1, &imgOutBarrier); - VK_CHECK_CALL(vkEndCommandBuffer(vk_computeCmdBuffer)); - - VkSubmitInfo submission = {}; - submission.sType = VK_STRUCTURE_TYPE_SUBMIT_INFO; - submission.commandBufferCount = 1; - submission.pCommandBuffers = &vk_computeCmdBuffer; - - VK_CHECK_CALL(vkQueueSubmit(vk_compute_queue, 1 /*submitCount*/, - &submission, VK_NULL_HANDLE /*fence*/)); - VK_CHECK_CALL(vkQueueWaitIdle(vk_compute_queue)); + submitCommandBuffer(vkCtx, commandBuffer, pool); } // Allocate temporary staging buffer and copy input data to device. - printString("Allocating staging memory and copying to device image\n"); + std::cout << "Allocating staging memory and copying to device image\n"; { - VkBuffer stagingBuffer; - VkDeviceMemory stagingMemory; - - stagingBuffer = vkutil::createBuffer(imgSizeBytes, - VK_BUFFER_USAGE_TRANSFER_SRC_BIT | - VK_BUFFER_USAGE_TRANSFER_DST_BIT); - auto inputStagingMemoryTypeIndex = vkutil::getBufferMemoryTypeIndex( - stagingBuffer, VK_MEMORY_PROPERTY_HOST_VISIBLE_BIT | - VK_MEMORY_PROPERTY_HOST_COHERENT_BIT); - stagingMemory = - vkutil::allocateDeviceMemory(imgSizeBytes, inputStagingMemoryTypeIndex, - nullptr /*image*/, false /*exportable*/); - VK_CHECK_CALL(vkBindBufferMemory(vk_device, stagingBuffer, stagingMemory, - 0 /*memoryOffset*/)); + auto staging = createStagingBuffer(vkCtx, imgSizeBytes, + VK_BUFFER_USAGE_TRANSFER_SRC_BIT | + VK_BUFFER_USAGE_TRANSFER_DST_BIT); // Copy host data to temporary staging buffer. float *inputStagingData = nullptr; - VK_CHECK_CALL(vkMapMemory(vk_device, stagingMemory, 0 /*offset*/, - imgSizeBytes, 0 /*flags*/, - (void **)&inputStagingData)); + VK_CHECK(vkMapMemory(vkCtx.device, staging.memory, 0 /*offset*/, + imgSizeBytes, 0 /*flags*/, + (void **)&inputStagingData)); for (int i = 0; i < (imgSizeElems); ++i) { inputStagingData[i] = inputVec[i]; } - vkUnmapMemory(vk_device, stagingMemory); + vkUnmapMemory(vkCtx.device, staging.memory); // Copy temporary staging buffer to device image memory. VkCommandBufferBeginInfo cbbi = {}; @@ -243,62 +209,40 @@ bool runTest(const sycl::device &syclDevice, sycl::range<2> dims, copyRegion.imageSubresource.aspectMask = VK_IMAGE_ASPECT_DEPTH_BIT; copyRegion.imageSubresource.layerCount = 1; - VK_CHECK_CALL(vkBeginCommandBuffer(vk_transferCmdBuffers[0], &cbbi)); - vkCmdCopyBufferToImage(vk_transferCmdBuffers[0], stagingBuffer, - vkInputImage, VK_IMAGE_LAYOUT_GENERAL, - 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)); + VkCommandPool pool; + VkCommandBuffer commandBuffer = createCommandBuffer(vkCtx, pool); + VK_CHECK(vkBeginCommandBuffer(commandBuffer, &cbbi)); + vkCmdCopyBufferToImage(commandBuffer, staging.buffer, inputImage.image, + VK_IMAGE_LAYOUT_GENERAL, 1 /*regionCount*/, + ©Region); + submitCommandBuffer(vkCtx, commandBuffer, pool); // Destroy temporary staging buffer and free memory. - vkDestroyBuffer(vk_device, stagingBuffer, nullptr); - vkFreeMemory(vk_device, stagingMemory, nullptr); + cleanupBuffer(vkCtx, staging); } - printString("Getting memory interop handles\n"); + std::cout << "Getting memory interop handles\n"; // Get memory interop handles. #ifdef _WIN32 - auto imgMemIn = vkutil::getMemoryWin32Handle(vkInputImageMemory); - auto imgMemOut = vkutil::getMemoryWin32Handle(vkOutputImageMemory); + auto imgMemIn = getMemHandle(vkCtx, inputImage.memory); + auto imgMemOut = getMemHandle(vkCtx, outputImage.memory); #else - auto imgMemIn = vkutil::getMemoryOpaqueFD(vkInputImageMemory); - auto imgMemOut = vkutil::getMemoryOpaqueFD(vkOutputImageMemory); + auto imgMemIn = getMemFd(vkCtx, inputImage.memory); + auto imgMemOut = getMemFd(vkCtx, outputImage.memory); #endif // Call into SYCL to fetch from input image, and populate the output image. - printString("Calling into SYCL with interop memory handles\n"); + std::cout << "Calling into SYCL with interop memory handles\n"; // Pass the real import size so the SYCL import matches the Vulkan allocation. runSycl(syclDevice, dims, localSize, imgMemIn, imgMemOut, importSizeBytes); // Copy image memory to temporary staging buffer, and back to host. - printString("Copying image memory to host\n"); + std::cout << "Copying image memory to host\n"; std::vector outputVec(imgSizeElems, 0.f); { - VkBuffer stagingBuffer; - VkDeviceMemory stagingMemory; - - stagingBuffer = vkutil::createBuffer(imgSizeBytes, - 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(imgSizeBytes, outputStagingMemoryTypeIndex, - nullptr /*image*/, false /*exportable*/); - VK_CHECK_CALL(vkBindBufferMemory(vk_device, stagingBuffer, stagingMemory, - 0 /*memoryOffset*/)); + auto staging = createStagingBuffer(vkCtx, imgSizeBytes, + VK_BUFFER_USAGE_TRANSFER_SRC_BIT | + VK_BUFFER_USAGE_TRANSFER_DST_BIT); VkCommandBufferBeginInfo cbbi = {}; cbbi.sType = VK_STRUCTURE_TYPE_COMMAND_BUFFER_BEGIN_INFO; @@ -309,44 +253,31 @@ bool runTest(const sycl::device &syclDevice, sycl::range<2> dims, copyRegion.imageSubresource.aspectMask = VK_IMAGE_ASPECT_DEPTH_BIT; copyRegion.imageSubresource.layerCount = 1; - VK_CHECK_CALL(vkBeginCommandBuffer(vk_transferCmdBuffers[1], &cbbi)); - vkCmdCopyImageToBuffer(vk_transferCmdBuffers[1], vkOutputImage, - VK_IMAGE_LAYOUT_GENERAL, stagingBuffer, + VkCommandPool pool; + VkCommandBuffer commandBuffer = createCommandBuffer(vkCtx, pool); + VK_CHECK(vkBeginCommandBuffer(commandBuffer, &cbbi)); + vkCmdCopyImageToBuffer(commandBuffer, outputImage.image, + VK_IMAGE_LAYOUT_GENERAL, staging.buffer, 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)); + submitCommandBuffer(vkCtx, commandBuffer, pool); // Copy temporary staging buffer output data to host output vector. float *outputStagingData = (float *)outputVec.data(); - VK_CHECK_CALL(vkMapMemory(vk_device, stagingMemory, 0 /*offset*/, - imgSizeBytes, 0 /*flags*/, - (void **)&outputStagingData)); + VK_CHECK(vkMapMemory(vkCtx.device, staging.memory, 0 /*offset*/, + imgSizeBytes, 0 /*flags*/, + (void **)&outputStagingData)); for (int i = 0; i < (imgSizeElems); ++i) { outputVec[i] = outputStagingData[i]; } - vkUnmapMemory(vk_device, stagingMemory); + vkUnmapMemory(vkCtx.device, staging.memory); // Destroy temporary staging buffer and free memory. - vkDestroyBuffer(vk_device, stagingBuffer, nullptr); - vkFreeMemory(vk_device, stagingMemory, nullptr); + cleanupBuffer(vkCtx, staging); } // Destroy images and free their memory. - vkDestroyImage(vk_device, vkInputImage, nullptr); - vkDestroyImage(vk_device, vkOutputImage, nullptr); - vkFreeMemory(vk_device, vkInputImageMemory, nullptr); - vkFreeMemory(vk_device, vkOutputImageMemory, nullptr); + cleanupImageResources(vkCtx, inputImage); + cleanupImageResources(vkCtx, outputImage); // Validate that SYCL made changes to the memory. bool validated = true; @@ -364,43 +295,32 @@ bool runTest(const sycl::device &syclDevice, sycl::range<2> dims, } if (validated) { - printString("Results are correct!\n"); + std::cout << "Results are correct!\n"; } return validated; } int main() { + try { + sycl::device syclDevice; + VulkanContext vkCtx = createSyclVulkanContext(syclDevice); + struct VulkanContextGuard { + VulkanContext &context; + ~VulkanContextGuard() { cleanupVulkanContext(context); } + } guard{vkCtx}; + + auto testPassed = runTest(vkCtx, syclDevice, {128, 128}, {16, 16}); + + if (testPassed) { + std::cout << "Test passed!\n"; + return EXIT_SUCCESS; + } - if (vkutil::setupInstance() != VK_SUCCESS) { - std::cerr << "Instance setup failed!\n"; - return EXIT_FAILURE; - } - - sycl::device syclDevice; - - if (vkutil::setupDevice(syclDevice) != VK_SUCCESS) { - std::cerr << "Device setup failed!\n"; - return EXIT_FAILURE; - } - - if (vkutil::setupCommandBuffers() != VK_SUCCESS) { - std::cerr << "Command buffers setup failed!\n"; + std::cerr << "Test failed\n"; return EXIT_FAILURE; - } - - auto testPassed = runTest(syclDevice, {128, 128}, {16, 16}); - - if (vkutil::cleanup() != VK_SUCCESS) { - std::cerr << "Cleanup failed!\n"; + } catch (const std::exception &e) { + std::cerr << "Vulkan interop test failed: " << e.what() << "\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/external_semaphore_regular_cl_fails.cpp b/sycl/test-e2e/bindless_images/vulkan_interop/external_semaphore_regular_cl_fails.cpp index c7616be3060db..d33bb10c8bab2 100644 --- a/sycl/test-e2e/bindless_images/vulkan_interop/external_semaphore_regular_cl_fails.cpp +++ b/sycl/test-e2e/bindless_images/vulkan_interop/external_semaphore_regular_cl_fails.cpp @@ -22,7 +22,7 @@ // that explicitly opts into no_immediate_command_list, and // expecting a sycl::exception. -#include "vulkan_setup.hpp" +#include "sycl_vulkan_setup.hpp" #include #include #include @@ -31,7 +31,7 @@ namespace syclexp = sycl::ext::oneapi::experimental; int main() { - VulkanContext vkCtx = createVulkanContext(); + VulkanContext vkCtx = createSyclVulkanContext(); VkSemaphore vkSem = createExportableSemaphore(vkCtx); // Lawful queue: import the semaphore here. diff --git a/sycl/test-e2e/bindless_images/vulkan_interop/mipmaps.cpp b/sycl/test-e2e/bindless_images/vulkan_interop/mipmaps.cpp index 62c4517bc4fe8..ee602d8121120 100644 --- a/sycl/test-e2e/bindless_images/vulkan_interop/mipmaps.cpp +++ b/sycl/test-e2e/bindless_images/vulkan_interop/mipmaps.cpp @@ -13,8 +13,8 @@ #define NOMINMAX #include -#include "../../CommonUtils/vulkan_common.hpp" #include "../helpers/common.hpp" +#include "sycl_vulkan_setup.hpp" #include #include @@ -63,10 +63,10 @@ handles_t create_handles(sycl::context &ctxt, sycl::device &dev, template -bool run_sycl(sycl::range globalSize, sycl::range localSize, +bool run_sycl(const sycl::device &dev, sycl::range globalSize, + sycl::range localSize, InteropMemHandleT inputImgInteropHandle, size_t mipLevels, size_t reqSize) { - sycl::device dev; sycl::queue q(dev); auto ctxt = q.get_context(); @@ -156,7 +156,7 @@ bool run_sycl(sycl::range globalSize, sycl::range localSize, 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) { + } catch (const sycl::exception &e) { std::cerr << "\tKernel submission failed! " << e.what() << std::endl; exit(-1); } catch (...) { @@ -164,7 +164,7 @@ bool run_sycl(sycl::range globalSize, sycl::range localSize, exit(-1); } - printString("Validating\n"); + std::cout << "Validating\n"; // Expected is sum of first two levels in the mipmap // Each subsequent level repeats in each dimension bool validated = true; @@ -233,7 +233,7 @@ bool run_sycl(sycl::range globalSize, sycl::range localSize, } } if (validated) { - printString("Results are correct!\n"); + std::cout << "Results are correct!\n"; } return validated; @@ -242,7 +242,8 @@ bool run_sycl(sycl::range globalSize, sycl::range localSize, template -bool run_test(sycl::range dims, sycl::range localSize, +bool run_test(VulkanContext &vkCtx, sycl::range dims, + const sycl::device &dev, sycl::range localSize, size_t mipLevels, unsigned int seed = 0) { uint32_t width = static_cast(dims[0]); @@ -264,42 +265,30 @@ bool run_test(sycl::range dims, sycl::range localSize, } using VecType = sycl::vec; - VkFormat format = vkutil::to_vulkan_format(COrder, CType); + VkFormat format = getVulkanFormat(NChannels); - printString("Creating input image\n"); + std::cout << "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, - mipLevels); + auto inputImage = createExportableImage( + vkCtx, {width, height, depth}, format, imgType, VK_IMAGE_TILING_OPTIMAL, + VK_IMAGE_USAGE_TRANSFER_SRC_BIT | VK_IMAGE_USAGE_TRANSFER_DST_BIT, + mipLevels); VkMemoryRequirements memRequirements; - auto inputImageMemoryTypeIndex = vkutil::getImageMemoryTypeIndex( - inputImage, VK_MEMORY_PROPERTY_DEVICE_LOCAL_BIT, memRequirements); - auto inputMemory = vkutil::allocateDeviceMemory( - memRequirements.size, inputImageMemoryTypeIndex, inputImage); - VK_CHECK_CALL(vkBindImageMemory(vk_device, inputImage, inputMemory, - 0 /*memoryOffset*/)); - - printString("Creating staging buffers\n"); + vkGetImageMemoryRequirements(vkCtx.device, inputImage.image, + &memRequirements); + + std::cout << "Creating staging buffers\n"; // Create input staging memory - auto inputStagingBuffer = vkutil::createBuffer( - memRequirements.size, - VK_BUFFER_USAGE_TRANSFER_SRC_BIT | VK_BUFFER_USAGE_TRANSFER_DST_BIT); - auto inputStagingMemoryTypeIndex = vkutil::getBufferMemoryTypeIndex( - inputStagingBuffer, VK_MEMORY_PROPERTY_HOST_VISIBLE_BIT | - VK_MEMORY_PROPERTY_HOST_COHERENT_BIT); - auto inputStagingMemory = vkutil::allocateDeviceMemory( - memRequirements.size, inputStagingMemoryTypeIndex, nullptr /*image*/, - false /*exportable*/); - VK_CHECK_CALL(vkBindBufferMemory(vk_device, inputStagingBuffer, - inputStagingMemory, 0 /*memoryOffset*/)); - - printString("Populating staging buffer\n"); + auto inputStaging = createStagingBuffer(vkCtx, memRequirements.size, + VK_BUFFER_USAGE_TRANSFER_SRC_BIT | + VK_BUFFER_USAGE_TRANSFER_DST_BIT); + + std::cout << "Populating staging buffer\n"; // Populate staging memory VecType *inputStagingData = nullptr; - VK_CHECK_CALL(vkMapMemory(vk_device, inputStagingMemory, 0 /*offset*/, - memRequirements.size, 0 /*flags*/, - (void **)&inputStagingData)); + VK_CHECK(vkMapMemory(vkCtx.device, inputStaging.memory, 0 /*offset*/, + memRequirements.size, 0 /*flags*/, + (void **)&inputStagingData)); // Set input data as each mip level -- 0 -> mip size e.g. (0,1,...,63,0,1,...) size_t offset = 0; @@ -314,35 +303,28 @@ bool run_test(sycl::range dims, sycl::range localSize, } offset += mipElems; } - vkUnmapMemory(vk_device, inputStagingMemory); + vkUnmapMemory(vkCtx.device, inputStaging.memory); - printString("Submitting image layout transition\n"); + std::cout << "Submitting image layout transition\n"; // Transition image layouts { VkImageMemoryBarrier barrierInput = - vkutil::createImageMemoryBarrier(inputImage, mipLevels); + createImageMemoryBarrier(inputImage.image, mipLevels); VkCommandBufferBeginInfo cbbi = {}; cbbi.sType = VK_STRUCTURE_TYPE_COMMAND_BUFFER_BEGIN_INFO; cbbi.flags = VK_COMMAND_BUFFER_USAGE_ONE_TIME_SUBMIT_BIT; - VK_CHECK_CALL(vkBeginCommandBuffer(vk_computeCmdBuffer, &cbbi)); - vkCmdPipelineBarrier(vk_computeCmdBuffer, VK_PIPELINE_STAGE_TOP_OF_PIPE_BIT, + VkCommandPool pool; + VkCommandBuffer commandBuffer = createCommandBuffer(vkCtx, pool); + VK_CHECK(vkBeginCommandBuffer(commandBuffer, &cbbi)); + vkCmdPipelineBarrier(commandBuffer, VK_PIPELINE_STAGE_TOP_OF_PIPE_BIT, VK_PIPELINE_STAGE_TRANSFER_BIT, 0, 0, nullptr, 0, nullptr, 1, &barrierInput); - VK_CHECK_CALL(vkEndCommandBuffer(vk_computeCmdBuffer)); - - VkSubmitInfo submission = {}; - submission.sType = VK_STRUCTURE_TYPE_SUBMIT_INFO; - submission.commandBufferCount = 1; - submission.pCommandBuffers = &vk_computeCmdBuffer; - - VK_CHECK_CALL(vkQueueSubmit(vk_compute_queue, 1 /*submitCount*/, - &submission, VK_NULL_HANDLE /*fence*/)); - VK_CHECK_CALL(vkQueueWaitIdle(vk_compute_queue)); + submitCommandBuffer(vkCtx, commandBuffer, pool); } - printString("Copying staging memory to images\n"); + std::cout << "Copying staging memory to images\n"; // Copy staging to main image memory { VkDeviceSize currentOffset{0}; @@ -367,110 +349,90 @@ bool run_test(sycl::range dims, sycl::range localSize, std::max(depth >> i, (uint32_t)1) * NChannels * sizeof(DType); - VK_CHECK_CALL(vkBeginCommandBuffer(vk_transferCmdBuffers[0], &cbbi)); - vkCmdCopyBufferToImage(vk_transferCmdBuffers[0], inputStagingBuffer, - inputImage, VK_IMAGE_LAYOUT_GENERAL, + VkCommandPool pool; + VkCommandBuffer commandBuffer = createCommandBuffer(vkCtx, pool); + VK_CHECK(vkBeginCommandBuffer(commandBuffer, &cbbi)); + vkCmdCopyBufferToImage(commandBuffer, inputStaging.buffer, + inputImage.image, VK_IMAGE_LAYOUT_GENERAL, 1 /*regionCount*/, ©Region); - VK_CHECK_CALL(vkEndCommandBuffer(vk_transferCmdBuffers[0])); - - VkSubmitInfo submission = {}; - submission.sType = VK_STRUCTURE_TYPE_SUBMIT_INFO; - submission.commandBufferCount = 1; - submission.pCommandBuffers = &vk_transferCmdBuffers[0]; - - VK_CHECK_CALL(vkQueueSubmit(vk_transfer_queue, 1 /*submitCount*/, - &submission, VK_NULL_HANDLE /*fence*/)); - VK_CHECK_CALL(vkQueueWaitIdle(vk_transfer_queue)); + submitCommandBuffer(vkCtx, commandBuffer, pool); } } - printString("Getting memory file descriptors and calling into SYCL\n"); + std::cout << "Getting memory file descriptors and calling into SYCL\n"; // Pass memory to SYCL for modification #ifdef _WIN32 - auto inputMemHandle = vkutil::getMemoryWin32Handle(inputMemory); + auto inputMemHandle = getMemHandle(vkCtx, inputImage.memory); #else - auto inputMemHandle = vkutil::getMemoryOpaqueFD(inputMemory); + auto inputMemHandle = getMemFd(vkCtx, inputImage.memory); #endif bool result = run_sycl( - dims, localSize, inputMemHandle, mipLevels, memRequirements.size); + dev, dims, localSize, inputMemHandle, mipLevels, memRequirements.size); // Cleanup - vkDestroyBuffer(vk_device, inputStagingBuffer, nullptr); - vkDestroyImage(vk_device, inputImage, nullptr); - vkFreeMemory(vk_device, inputStagingMemory, nullptr); - vkFreeMemory(vk_device, inputMemory, nullptr); + cleanupBuffer(vkCtx, inputStaging); + cleanupImageResources(vkCtx, inputImage); return result; } -bool run_tests() { +bool run_tests(VulkanContext &vkCtx, const sycl::device &dev) { bool valid = run_test<2, float, 4, sycl::image_channel_type::fp32, sycl::image_channel_order::rgba, class float_2d>( - {16, 16}, {2, 2}, 2, 0); + vkCtx, dev, {16, 16}, {2, 2}, 2, 0); valid &= run_test<2, float, 2, sycl::image_channel_type::fp32, sycl::image_channel_order::rg, class float_2d_large>( - {8, 8}, {4, 2}, 2, 0); + vkCtx, dev, {8, 8}, {4, 2}, 2, 0); - valid &= run_test<3, char, 2, sycl::image_channel_type::signed_int8, + valid &= run_test<3, int8_t, 2, sycl::image_channel_type::signed_int8, sycl::image_channel_order::rg, class float_3d>( - {8, 8, 8}, {2, 2, 2}, 2, 0); + vkCtx, dev, {8, 8, 8}, {2, 2, 2}, 2, 0); valid &= run_test<2, uint32_t, 1, sycl::image_channel_type::unsigned_int32, sycl::image_channel_order::r, class uint32_2d>( - {32, 32}, {4, 2}, 2, 0); + vkCtx, dev, {32, 32}, {4, 2}, 2, 0); valid &= run_test<3, uint32_t, 4, sycl::image_channel_type::unsigned_int32, sycl::image_channel_order::rgba, class uint_3d_large>( - {8, 8, 8}, {2, 2, 4}, 2, 0); + vkCtx, dev, {8, 8, 8}, {2, 2, 4}, 2, 0); valid &= run_test<2, int32_t, 1, sycl::image_channel_type::signed_int32, - sycl::image_channel_order::r, class int32_2d>({64, 64}, - {4, 2}, 2, 0); + sycl::image_channel_order::r, class int32_2d>( + vkCtx, dev, {64, 64}, {4, 2}, 2, 0); valid &= run_test<3, int32_t, 2, sycl::image_channel_type::signed_int32, sycl::image_channel_order::rg, class int32_3d>( - {8, 8, 8}, {4, 2, 4}, 2, 0); + vkCtx, dev, {8, 8, 8}, {4, 2, 4}, 2, 0); valid &= run_test<3, int16_t, 1, sycl::image_channel_type::signed_int16, sycl::image_channel_order::r, class int16_3d>( - {32, 32, 32}, {4, 2, 4}, 2, 0); + vkCtx, dev, {32, 32, 32}, {4, 2, 4}, 2, 0); return valid; } int main() { + try { + sycl::device dev; + VulkanContext vkCtx = createSyclVulkanContext(dev); + struct VulkanContextGuard { + VulkanContext &context; + ~VulkanContextGuard() { cleanupVulkanContext(context); } + } guard{vkCtx}; + + bool result_ok = run_tests(vkCtx, dev); + + if (result_ok) { + std::cout << "All tests passed!\n"; + return EXIT_SUCCESS; + } - if (vkutil::setupInstance() != VK_SUCCESS) { - std::cerr << "Instance setup failed!\n"; - return EXIT_FAILURE; - } - - sycl::device dev; - - if (vkutil::setupDevice(dev) != VK_SUCCESS) { - std::cerr << "Device setup failed!\n"; - return EXIT_FAILURE; - } - - if (vkutil::setupCommandBuffers() != VK_SUCCESS) { - std::cerr << "Compute pipeline setup failed!\n"; + std::cerr << "Test failed\n"; return EXIT_FAILURE; - } - - bool result_ok = run_tests(); - - if (vkutil::cleanup() != VK_SUCCESS) { - std::cerr << "Cleanup failed!\n"; + } catch (const std::exception &e) { + std::cerr << "Vulkan interop test failed: " << e.what() << "\n"; return EXIT_FAILURE; } - - if (result_ok) { - std::cout << "All tests passed!\n"; - return EXIT_SUCCESS; - } - - std::cerr << "Test failed\n"; - return EXIT_FAILURE; } diff --git a/sycl/test-e2e/bindless_images/vulkan_interop/sycl_vulkan_setup.hpp b/sycl/test-e2e/bindless_images/vulkan_interop/sycl_vulkan_setup.hpp new file mode 100644 index 0000000000000..de13e441fe6b3 --- /dev/null +++ b/sycl/test-e2e/bindless_images/vulkan_interop/sycl_vulkan_setup.hpp @@ -0,0 +1,21 @@ +#pragma once + +#include + +#include +#include +#include + +#include "vulkan_setup.hpp" + +inline VulkanContext createSyclVulkanContext(const sycl::device &SyclDevice) { + if (!SyclDevice.has(sycl::aspect::ext_intel_device_info_uuid)) + throw std::runtime_error("SYCL device UUID is unavailable!"); + + return createVulkanContext( + SyclDevice.get_info()); +} + +inline VulkanContext createSyclVulkanContext() { + return createSyclVulkanContext(sycl::device{}); +} diff --git a/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_setup.hpp b/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_setup.hpp index d7b5a0e409f58..c6d1ba0577c65 100644 --- a/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_setup.hpp +++ b/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_setup.hpp @@ -1,5 +1,6 @@ #pragma once +#include #include #include #include @@ -302,6 +303,19 @@ template inline bool checkValue(T actual, T expected) { } } +namespace util { + +template +bool is_equal(DType lhs, DType rhs, float epsilon = 0.0001f) { + if constexpr (std::is_floating_point_v) { + return std::abs(lhs - rhs) < epsilon; + } else { + return lhs == rhs; + } +} + +} // namespace util + // --------------------------------------------------------- // Boilerplate // --------------------------------------------------------- @@ -336,7 +350,8 @@ inline uint32_t findMemoryType(VkPhysicalDevice physicalDevice, throw std::runtime_error("failed to find suitable memory type!"); } -inline VulkanContext createVulkanContext() { +inline VulkanContext +createVulkanContext(const std::array &SyclDeviceUUID) { VulkanContext ctx; VkApplicationInfo appInfo{}; appInfo.sType = VK_STRUCTURE_TYPE_APPLICATION_INFO; @@ -434,7 +449,24 @@ inline VulkanContext createVulkanContext() { vkEnumeratePhysicalDevices(ctx.instance, &deviceCount, nullptr); std::vector devices(deviceCount); vkEnumeratePhysicalDevices(ctx.instance, &deviceCount, devices.data()); - ctx.physicalDevice = devices[0]; + + ctx.physicalDevice = VK_NULL_HANDLE; + for (VkPhysicalDevice device : devices) { + VkPhysicalDeviceIDProperties idProperties{}; + idProperties.sType = VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_ID_PROPERTIES; + VkPhysicalDeviceProperties2 properties{}; + properties.sType = VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_PROPERTIES_2; + properties.pNext = &idProperties; + vkGetPhysicalDeviceProperties2(device, &properties); + if (std::memcmp(idProperties.deviceUUID, SyclDeviceUUID.data(), + VK_UUID_SIZE) == 0) { + ctx.physicalDevice = device; + break; + } + } + + if (ctx.physicalDevice == VK_NULL_HANDLE) + throw std::runtime_error("Failed to find matching Vulkan physical device!"); uint32_t queueFamilyCount = 0; vkGetPhysicalDeviceQueueFamilyProperties(ctx.physicalDevice, @@ -495,18 +527,18 @@ inline void cleanupVulkanContext(VulkanContext &ctx) { #endif vkDestroyInstance(ctx.instance, nullptr); } - inline ImageResources createExportableImage( VulkanContext &ctx, VkExtent3D extent, VkFormat format, VkImageType type, VkImageTiling tiling = VK_IMAGE_TILING_OPTIMAL, VkImageUsageFlags usage = VK_IMAGE_USAGE_STORAGE_BIT | VK_IMAGE_USAGE_SAMPLED_BIT | VK_IMAGE_USAGE_TRANSFER_SRC_BIT | - VK_IMAGE_USAGE_TRANSFER_DST_BIT) { + VK_IMAGE_USAGE_TRANSFER_DST_BIT, + uint32_t mipLevels = 1) { VkImageCreateInfo imageInfo = {VK_STRUCTURE_TYPE_IMAGE_CREATE_INFO}; imageInfo.imageType = type; imageInfo.extent = extent; - imageInfo.mipLevels = 1; + imageInfo.mipLevels = mipLevels; imageInfo.arrayLayers = 1; imageInfo.format = format; imageInfo.tiling = tiling; @@ -672,6 +704,57 @@ inline void cleanupBuffer(VulkanContext &ctx, BufferResources &res) { vkFreeMemory(ctx.device, res.memory, nullptr); } +inline VkCommandBuffer createCommandBuffer(const VulkanContext &ctx, + VkCommandPool &pool) { + VkCommandPoolCreateInfo poolInfo{}; + poolInfo.sType = VK_STRUCTURE_TYPE_COMMAND_POOL_CREATE_INFO; + poolInfo.queueFamilyIndex = ctx.queueFamilyIndex; + poolInfo.flags = VK_COMMAND_POOL_CREATE_RESET_COMMAND_BUFFER_BIT; + VK_CHECK(vkCreateCommandPool(ctx.device, &poolInfo, nullptr, &pool)); + + VkCommandBufferAllocateInfo allocInfo{}; + allocInfo.sType = VK_STRUCTURE_TYPE_COMMAND_BUFFER_ALLOCATE_INFO; + allocInfo.commandPool = pool; + allocInfo.level = VK_COMMAND_BUFFER_LEVEL_PRIMARY; + allocInfo.commandBufferCount = 1; + VkCommandBuffer commandBuffer; + VK_CHECK(vkAllocateCommandBuffers(ctx.device, &allocInfo, &commandBuffer)); + return commandBuffer; +} + +inline void submitCommandBuffer(VulkanContext &ctx, + VkCommandBuffer commandBuffer, + VkCommandPool pool) { + VK_CHECK(vkEndCommandBuffer(commandBuffer)); + VkSubmitInfo submitInfo{VK_STRUCTURE_TYPE_SUBMIT_INFO}; + submitInfo.commandBufferCount = 1; + submitInfo.pCommandBuffers = &commandBuffer; + VK_CHECK(vkQueueSubmit(ctx.queue, 1, &submitInfo, VK_NULL_HANDLE)); + VK_CHECK(vkQueueWaitIdle(ctx.queue)); + vkDestroyCommandPool(ctx.device, pool, nullptr); +} + +inline VkImageMemoryBarrier createImageMemoryBarrier( + VkImage image, uint32_t mipLevels, + VkImageLayout oldLayout = VK_IMAGE_LAYOUT_UNDEFINED, + VkImageLayout newLayout = VK_IMAGE_LAYOUT_GENERAL, + VkAccessFlags srcAccessMask = 0, + VkAccessFlags dstAccessMask = VK_ACCESS_TRANSFER_WRITE_BIT) { + VkImageMemoryBarrier barrier{}; + barrier.sType = VK_STRUCTURE_TYPE_IMAGE_MEMORY_BARRIER; + barrier.srcAccessMask = srcAccessMask; + barrier.dstAccessMask = dstAccessMask; + barrier.oldLayout = oldLayout; + barrier.newLayout = newLayout; + barrier.srcQueueFamilyIndex = VK_QUEUE_FAMILY_IGNORED; + barrier.dstQueueFamilyIndex = VK_QUEUE_FAMILY_IGNORED; + barrier.image = image; + barrier.subresourceRange.aspectMask = VK_IMAGE_ASPECT_COLOR_BIT; + barrier.subresourceRange.levelCount = mipLevels; + barrier.subresourceRange.layerCount = 1; + return barrier; +} + // --------------------------------------------------------- // PLATFORM SPECIFIC GETTERS // --------------------------------------------------------- diff --git a/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_2d_arithmetic.cpp b/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_2d_arithmetic.cpp index bdecc738b7467..84f224588dab0 100644 --- a/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_2d_arithmetic.cpp +++ b/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_2d_arithmetic.cpp @@ -115,7 +115,7 @@ // clang-format on #include -#include "vulkan_setup.hpp" +#include "sycl_vulkan_setup.hpp" #include #include @@ -208,7 +208,7 @@ int runTest( std::cout << " Mode: " << (useSampled ? "SAMPLED Input" : "UNSAMPLED Input") << std::endl; - VulkanContext vkCtx = createVulkanContext(); + VulkanContext vkCtx = createSyclVulkanContext(); VkExtent3D extent = {(uint32_t)width, (uint32_t)height, 1}; VkImageUsageFlags usage = VK_IMAGE_USAGE_TRANSFER_SRC_BIT | diff --git a/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_buffer.cpp b/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_buffer.cpp index 94aae568953a5..6dcd2f21863a0 100644 --- a/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_buffer.cpp +++ b/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_buffer.cpp @@ -46,7 +46,7 @@ // clang-format on #include -#include "vulkan_setup.hpp" +#include "sycl_vulkan_setup.hpp" #include #include #include @@ -234,7 +234,7 @@ int main(int argc, char **argv) { << " | Type: " << (useDmaBuf ? "DMA_BUF" : "OPAQUE") << std::endl; // VULKAN SETUP - VulkanContext vkCtx = createVulkanContext(); + VulkanContext vkCtx = createSyclVulkanContext(); VkDeviceSize bufferSize = numElements * sizeof(uint32_t); // Create Input and Output Buffers diff --git a/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_buffer_binary_semaphore.cpp b/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_buffer_binary_semaphore.cpp index e167ed8b01b6c..541811975aa06 100644 --- a/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_buffer_binary_semaphore.cpp +++ b/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_buffer_binary_semaphore.cpp @@ -41,7 +41,7 @@ */ // clang-format on -#include "vulkan_setup.hpp" +#include "sycl_vulkan_setup.hpp" #include #include #include @@ -84,7 +84,7 @@ int main(int argc, char **argv) { << " | Mode: " << modeStr << std::endl; // VULKAN SETUP - VulkanContext vkCtx = createVulkanContext(); + VulkanContext vkCtx = createSyclVulkanContext(); // Exportable device-local buffers BufferResources inBuf = createExportableBuffer( diff --git a/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_buffer_timeline_semaphore.cpp b/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_buffer_timeline_semaphore.cpp index 0f4cd28dd7091..018e4f0a29458 100644 --- a/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_buffer_timeline_semaphore.cpp +++ b/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_buffer_timeline_semaphore.cpp @@ -43,7 +43,7 @@ */ // clang-format on -#include "vulkan_setup.hpp" +#include "sycl_vulkan_setup.hpp" #include #include #include @@ -79,7 +79,7 @@ int main(int argc, char **argv) { << " | Semaphores: " << (useSemaphores ? "ON" : "OFF") << std::endl; // VULKAN SETUP - VulkanContext vkCtx = createVulkanContext(); + VulkanContext vkCtx = createSyclVulkanContext(); // Exportable device-local buffers BufferResources inBuf = createExportableBuffer( diff --git a/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_image_interop_read_1d.cpp b/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_image_interop_read_1d.cpp index d20c957c82c21..2c42b58421014 100644 --- a/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_image_interop_read_1d.cpp +++ b/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_image_interop_read_1d.cpp @@ -166,7 +166,7 @@ VK_FORMAT_R8G8B8A8_UNORM // clang-format on #include -#include "vulkan_setup.hpp" +#include "sycl_vulkan_setup.hpp" #include #include @@ -250,7 +250,7 @@ int runTest( std::cout << "VK Format: " << getFormatString(vkFormat) << std::endl; // Setup Vulkan - VulkanContext vkCtx = createVulkanContext(); + VulkanContext vkCtx = createSyclVulkanContext(); VkExtent3D extent = {(uint32_t)width, 1, 1}; ImageResources imgRes = createExportableImage(vkCtx, extent, vkFormat, VK_IMAGE_TYPE_1D, tiling); diff --git a/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_image_interop_read_2d.cpp b/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_image_interop_read_2d.cpp index 965f39fd859f5..b25a92378d7ec 100644 --- a/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_image_interop_read_2d.cpp +++ b/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_image_interop_read_2d.cpp @@ -121,7 +121,7 @@ // clang-format on #include -#include "vulkan_setup.hpp" +#include "sycl_vulkan_setup.hpp" #include #include @@ -222,7 +222,7 @@ int runTest( std::cout << "VK Format: " << getFormatString(vkFormat) << std::endl; // Setup Vulkan - VulkanContext vkCtx = createVulkanContext(); + VulkanContext vkCtx = createSyclVulkanContext(); VkExtent3D extent = {(uint32_t)width, (uint32_t)height, 1}; ImageResources imgRes = createExportableImage(vkCtx, extent, vkFormat, VK_IMAGE_TYPE_2D, tiling); diff --git a/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_image_interop_read_3d.cpp b/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_image_interop_read_3d.cpp index 6f62453eece48..f570d5ac9a90f 100644 --- a/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_image_interop_read_3d.cpp +++ b/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_image_interop_read_3d.cpp @@ -35,7 +35,7 @@ ./vsr_3d_test.bin --linear --type unorm8 16x16x16 */ // clang-format on -#include "vulkan_setup.hpp" +#include "sycl_vulkan_setup.hpp" #include #include @@ -121,7 +121,7 @@ int runTest( std::cout << "VK Format: " << getFormatString(vkFormat) << std::endl; // Setup Vulkan - VulkanContext vkCtx = createVulkanContext(); + VulkanContext vkCtx = createSyclVulkanContext(); VkExtent3D extent = {(uint32_t)width, (uint32_t)height, (uint32_t)depth}; ImageResources imgRes = createExportableImage(vkCtx, extent, vkFormat, VK_IMAGE_TYPE_3D, tiling); diff --git a/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_image_interop_write_1d_unsampled.cpp b/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_image_interop_write_1d_unsampled.cpp index 631950021ab96..0df0df0361847 100644 --- a/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_image_interop_write_1d_unsampled.cpp +++ b/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_image_interop_write_1d_unsampled.cpp @@ -93,7 +93,7 @@ // clang-format on #include -#include "vulkan_setup.hpp" +#include "sycl_vulkan_setup.hpp" #include #include @@ -165,7 +165,7 @@ int runTest( : getVulkanFormat(channels); std::cout << "VK Format: " << getFormatString(vkFormat) << std::endl; - VulkanContext vkCtx = createVulkanContext(); + VulkanContext vkCtx = createSyclVulkanContext(); VkExtent3D extent = {(uint32_t)width, 1, 1}; ImageResources imgRes = createExportableImage(vkCtx, extent, vkFormat, VK_IMAGE_TYPE_1D, tiling); diff --git a/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_image_interop_write_2d_unsampled.cpp b/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_image_interop_write_2d_unsampled.cpp index e79f68407a98a..5d8200aad1f72 100644 --- a/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_image_interop_write_2d_unsampled.cpp +++ b/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_image_interop_write_2d_unsampled.cpp @@ -81,7 +81,7 @@ // clang-format on #include -#include "vulkan_setup.hpp" +#include "sycl_vulkan_setup.hpp" #include #include @@ -152,7 +152,7 @@ int runTest( : getVulkanFormat(channels); std::cout << "VK Format: " << getFormatString(vkFormat) << std::endl; - VulkanContext vkCtx = createVulkanContext(); + VulkanContext vkCtx = createSyclVulkanContext(); VkExtent3D extent = {(uint32_t)width, (uint32_t)height, 1}; ImageResources imgRes = createExportableImage(vkCtx, extent, vkFormat, VK_IMAGE_TYPE_2D, tiling); diff --git a/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_image_interop_write_3d_unsampled.cpp b/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_image_interop_write_3d_unsampled.cpp index 381d0caa1f717..02d6944c2b591 100644 --- a/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_image_interop_write_3d_unsampled.cpp +++ b/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_image_interop_write_3d_unsampled.cpp @@ -32,7 +32,7 @@ // clang-format on #include -#include "vulkan_setup.hpp" +#include "sycl_vulkan_setup.hpp" #include #include @@ -103,7 +103,7 @@ int runTest( : getVulkanFormat(channels); std::cout << "VK Format: " << getFormatString(vkFormat) << std::endl; - VulkanContext vkCtx = createVulkanContext(); + VulkanContext vkCtx = createSyclVulkanContext(); VkExtent3D extent = {(uint32_t)width, (uint32_t)height, (uint32_t)depth}; ImageResources imgRes = createExportableImage(vkCtx, extent, vkFormat, VK_IMAGE_TYPE_3D, tiling); diff --git a/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_image_unsampled_timeline_semaphore.cpp b/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_image_unsampled_timeline_semaphore.cpp index fb13ce6323da2..98a8cc4d50463 100644 --- a/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_image_unsampled_timeline_semaphore.cpp +++ b/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_image_unsampled_timeline_semaphore.cpp @@ -48,7 +48,7 @@ // clang-format on #include -#include "vulkan_setup.hpp" +#include "sycl_vulkan_setup.hpp" #include #include #include @@ -109,7 +109,7 @@ int runTest( std::cout << " VK Format: " << getFormatString(vkFormat) << std::endl; - VulkanContext vkCtx = createVulkanContext(); + VulkanContext vkCtx = createSyclVulkanContext(); VkExtent3D extent = {(uint32_t)imageWidth, (uint32_t)imageHeight, 1}; VkImageUsageFlags usage = VK_IMAGE_USAGE_TRANSFER_SRC_BIT |