diff --git a/projects/hip-tests/catch/unit/CMakeLists.txt b/projects/hip-tests/catch/unit/CMakeLists.txt index d37cea450b..57eb632947 100644 --- a/projects/hip-tests/catch/unit/CMakeLists.txt +++ b/projects/hip-tests/catch/unit/CMakeLists.txt @@ -34,6 +34,8 @@ add_subdirectory(multiThread) add_subdirectory(compiler) add_subdirectory(errorHandling) add_subdirectory(cooperativeGrps) -#if(HIP_PLATFORM STREQUAL "amd") +if(HIP_PLATFORM STREQUAL "amd") #add_subdirectory(clock) -#endif() +# Vulkan interop APIs currently undefined for Nvidia +add_subdirectory(vulkan_interop) +endif() diff --git a/projects/hip-tests/catch/unit/vulkan_interop/CMakeLists.txt b/projects/hip-tests/catch/unit/vulkan_interop/CMakeLists.txt new file mode 100644 index 0000000000..2f015d726a --- /dev/null +++ b/projects/hip-tests/catch/unit/vulkan_interop/CMakeLists.txt @@ -0,0 +1,26 @@ +set(TEST_SRC + vulkan_test.cc + hipExternalMemoryGetMappedBuffer.cc + ) + +find_package(Vulkan) +if(NOT Vulkan_FOUND) + if(EXISTS "${VULKAN_PATH}") + message(STATUS "Vulkan SDK: ${VULKAN_PATH}") + elseif (EXISTS "$ENV{VULKAN_SDK}") + message(STATUS "FOUND VULKAN SDK: $ENV{VULKAN_SDK}") + set(VULKAN_PATH $ENV{VULKAN_SDK}) + else() + message("Error: Unable to locate Vulkan SDK. please specify VULKAN_PATH") + return() + endif() +endif() + +hip_add_exe_to_target(NAME VulkanInteropTest + TEST_SRC ${TEST_SRC} + TEST_TARGET_NAME build_tests) +if (WIN32) + target_link_libraries(VulkanInteropTest vulkan-1) +else (WIN32) + target_link_libraries(VulkanInteropTest vulkan) +endif (WIN32) \ No newline at end of file diff --git a/projects/hip-tests/catch/unit/vulkan_interop/hipExternalMemoryGetMappedBuffer.cc b/projects/hip-tests/catch/unit/vulkan_interop/hipExternalMemoryGetMappedBuffer.cc new file mode 100644 index 0000000000..c4ad1443c4 --- /dev/null +++ b/projects/hip-tests/catch/unit/vulkan_interop/hipExternalMemoryGetMappedBuffer.cc @@ -0,0 +1,153 @@ +/* +Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved. +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#include "vulkan_test.hh" + +constexpr bool enable_validation = false; + +template __global__ void Set(T* ptr, const T val) { ptr[threadIdx.x] = val; } + +TEST_CASE("Unit_hipExternalMemoryGetMappedBuffer_Vulkan_Positive_Read_Write") { + VulkanTest vkt(enable_validation); + using type = uint8_t; + constexpr uint32_t count = 3; + + const auto vk_storage = + vkt.CreateMappedStorage(count, VK_BUFFER_USAGE_TRANSFER_DST_BIT, true); + + const auto hip_ext_mem_desc = vkt.BuildMemoryDescriptor(vk_storage.memory, vk_storage.size); + hipExternalMemory_t hip_ext_memory; + HIP_CHECK(hipImportExternalMemory(&hip_ext_memory, &hip_ext_mem_desc)); + + hipExternalMemoryBufferDesc external_mem_buffer_desc = {}; + external_mem_buffer_desc.size = vk_storage.size; + + type* hip_dev_ptr = nullptr; + HIP_CHECK(hipExternalMemoryGetMappedBuffer(reinterpret_cast(&hip_dev_ptr), hip_ext_memory, + &external_mem_buffer_desc)); + REQUIRE(nullptr != hip_dev_ptr); + + vk_storage.host_ptr[0] = 41; + vk_storage.host_ptr[1] = 40; + vk_storage.host_ptr[2] = 43; + + std::vector read_buffer(count, 0); + HIP_CHECK( + hipMemcpy(read_buffer.data(), hip_dev_ptr, count * sizeof(type), hipMemcpyDeviceToHost)); + REQUIRE(41 == read_buffer[0]); + REQUIRE(40 == read_buffer[1]); + REQUIRE(43 == read_buffer[2]); + + Set<<<1, 1>>>(hip_dev_ptr + 1, static_cast(42)); + HIP_CHECK(hipDeviceSynchronize()); + REQUIRE(41 == vk_storage.host_ptr[0]); + REQUIRE(42 == vk_storage.host_ptr[1]); + REQUIRE(43 == vk_storage.host_ptr[2]); + + // Defect - EXSWHTEC-181 + // HIP_CHECK(hipFree(hip_dev_ptr)); + HIP_CHECK(hipDestroyExternalMemory(hip_ext_memory)); +} + +// Disabled on AMD due to defect - EXSWHTEC-175 +#if HT_NVIDIA +TEST_CASE("Unit_hipExternalMemoryGetMappedBuffer_Vulkan_Positive_Read_Write_With_Offset") { + VulkanTest vkt(enable_validation); + using type = uint8_t; + constexpr uint32_t count = 2; + + const auto vk_storage = + vkt.CreateMappedStorage(count, VK_BUFFER_USAGE_TRANSFER_DST_BIT, true); + + const auto hip_ext_mem_desc = vkt.BuildMemoryDescriptor(vk_storage.memory, vk_storage.size); + hipExternalMemory_t hip_ext_memory; + HIP_CHECK(hipImportExternalMemory(&hip_ext_memory, &hip_ext_mem_desc)); + + hipExternalMemoryBufferDesc external_mem_buffer_desc = {}; + constexpr auto offset = (count - 1) * sizeof(type); + external_mem_buffer_desc.size = vk_storage.size - offset; + external_mem_buffer_desc.offset = offset; + + type* hip_dev_ptr = nullptr; + HIP_CHECK(hipExternalMemoryGetMappedBuffer(reinterpret_cast(&hip_dev_ptr), hip_ext_memory, + &external_mem_buffer_desc)); + + vk_storage.host_ptr[0] = 41; + vk_storage.host_ptr[1] = 42; + type read_val = 0; + HIP_CHECK(hipMemcpy(&read_val, hip_dev_ptr, 1, hipMemcpyDeviceToHost)); + REQUIRE(42 == read_val); + + // Defect - EXSWHTEC-181 + // HIP_CHECK(hipFree(hip_dev_ptr)); + HIP_CHECK(hipDestroyExternalMemory(hip_ext_memory)); +} +#endif + +TEST_CASE("Unit_hipExternalMemoryGetMappedBuffer_Vulkan_Negative_Parameters") { + VulkanTest vkt(enable_validation); + const auto vk_storage = vkt.CreateMappedStorage(1, VK_BUFFER_USAGE_TRANSFER_DST_BIT, true); + + const auto hip_ext_mem_desc = vkt.BuildMemoryDescriptor(vk_storage.memory, vk_storage.size); + hipExternalMemory_t hip_ext_memory; + HIP_CHECK(hipImportExternalMemory(&hip_ext_memory, &hip_ext_mem_desc)); + + hipExternalMemoryBufferDesc external_mem_buffer_desc = {}; + external_mem_buffer_desc.size = vk_storage.size; + void* hip_dev_ptr = nullptr; + +// Disabled on AMD due to defect - EXSWHTEC-176 +#if HT_NVIDIA + SECTION("devPtr == nullptr") { + HIP_CHECK_ERROR( + hipExternalMemoryGetMappedBuffer(nullptr, hip_ext_memory, &external_mem_buffer_desc), + hipErrorInvalidValue); + } +#endif + +// Disabled on AMD due to defect - EXSWHTEC-177 +#if HT_NVIDIA + SECTION("bufferDesc == nullptr") { + HIP_CHECK_ERROR(hipExternalMemoryGetMappedBuffer(&hip_dev_ptr, hip_ext_memory, nullptr), + hipErrorInvalidValue); + } +#endif + +// Disabled on AMD due to defect - EXSWHTEC-179 +#if HT_NVIDIA + SECTION("bufferDesc.flags != 0") { + external_mem_buffer_desc.flags = 1; + HIP_CHECK_ERROR( + hipExternalMemoryGetMappedBuffer(&hip_dev_ptr, hip_ext_memory, &external_mem_buffer_desc), + hipErrorInvalidValue); + } +#endif + +// Disabled on AMD due to defect - EXSWHTEC-180 +#if HT_NVIDIA + SECTION("bufferDesc.offset + bufferDesc.size > hipExternalMemHandleDesc.size") { + external_mem_buffer_desc.offset = 1; + HIP_CHECK_ERROR( + hipExternalMemoryGetMappedBuffer(&hip_dev_ptr, hip_ext_memory, &external_mem_buffer_desc), + hipErrorInvalidValue); + } +#endif +} \ No newline at end of file diff --git a/projects/hip-tests/catch/unit/vulkan_interop/vulkan_test.cc b/projects/hip-tests/catch/unit/vulkan_interop/vulkan_test.cc new file mode 100644 index 0000000000..4252916177 --- /dev/null +++ b/projects/hip-tests/catch/unit/vulkan_interop/vulkan_test.cc @@ -0,0 +1,410 @@ +/* +Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved. +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#include "vulkan_test.hh" + +#include +#include + +VkFence VulkanTest::CreateFence() { + VkFence fence; + VkFenceCreateInfo fence_create_info = {}; + fence_create_info.sType = VK_STRUCTURE_TYPE_FENCE_CREATE_INFO; + fence_create_info.flags = 0; + VK_CHECK_RESULT(vkCreateFence(_device, &fence_create_info, nullptr, &fence)); + + _fences.push_back(fence); + return fence; +} + +VkSemaphore VulkanTest::CreateExternalSemaphore(VkSemaphoreType sem_type, uint64_t initial_value) { + VkExportSemaphoreCreateInfoKHR export_sem_create_info = {}; + export_sem_create_info.sType = VK_STRUCTURE_TYPE_EXPORT_SEMAPHORE_CREATE_INFO_KHR; + export_sem_create_info.handleTypes = _sem_handle_type; + + if (sem_type == VK_SEMAPHORE_TYPE_TIMELINE) { + VkSemaphoreTypeCreateInfo timeline_create_info = {}; + timeline_create_info.sType = VK_STRUCTURE_TYPE_SEMAPHORE_TYPE_CREATE_INFO; + timeline_create_info.semaphoreType = VK_SEMAPHORE_TYPE_TIMELINE; + timeline_create_info.initialValue = initial_value; + export_sem_create_info.pNext = &timeline_create_info; + } else { + export_sem_create_info.pNext = nullptr; + } + + VkSemaphoreCreateInfo semaphore_create_info = {}; + semaphore_create_info.sType = VK_STRUCTURE_TYPE_SEMAPHORE_CREATE_INFO; + semaphore_create_info.pNext = &export_sem_create_info; + + VkSemaphore semaphore; + VK_CHECK_RESULT(vkCreateSemaphore(_device, &semaphore_create_info, nullptr, &semaphore)); + + _semaphores.push_back(semaphore); + return semaphore; +} + +hipExternalSemaphoreHandleDesc VulkanTest::BuildSemaphoreDescriptor(VkSemaphore vk_sem, + VkSemaphoreType sem_type) { + hipExternalSemaphoreHandleDesc sem_handle_desc = {}; + sem_handle_desc.type = VulkanSemHandleTypeToHIPHandleType(sem_type); +#ifdef _WIN64 + sem_handle_desc.handle.win32.handle = GetSemaphoreHandle(vk_sem); +#else + sem_handle_desc.handle.fd = GetSemaphoreHandle(vk_sem); +#endif + sem_handle_desc.flags = 0; + + return sem_handle_desc; +} + +hipExternalMemoryHandleDesc VulkanTest::BuildMemoryDescriptor(VkDeviceMemory vk_mem, + uint32_t size) { + hipExternalMemoryHandleDesc mem_handle_desc = {}; + mem_handle_desc.type = VulkanMemHandleTypeToHIPHandleType(); +#ifdef _WIN64 + mem_handle_desc.handle.win32.handle = GetMemoryHandle(ck_mem); +#else + mem_handle_desc.handle.fd = GetMemoryHandle(vk_mem); +#endif + mem_handle_desc.size = size; + + return mem_handle_desc; +} + +void VulkanTest::CreateInstance() { + UNSCOPED_INFO("Not all of the required instance extensions are supported"); + REQUIRE(CheckExtensionSupport(_required_instance_extensions)); + if (_enable_validation) { + EnableValidationLayer(); + } + + VkApplicationInfo app_info = {}; + app_info.sType = VK_STRUCTURE_TYPE_APPLICATION_INFO; + app_info.apiVersion = VK_API_VERSION_1_2; + + VkInstanceCreateInfo create_info = {}; + create_info.sType = VK_STRUCTURE_TYPE_INSTANCE_CREATE_INFO; + create_info.pApplicationInfo = &app_info; + create_info.enabledExtensionCount = static_cast(_required_instance_extensions.size()); + create_info.ppEnabledExtensionNames = _required_instance_extensions.data(); + create_info.enabledLayerCount = static_cast(_enabled_layers.size()); + create_info.ppEnabledLayerNames = _enabled_layers.data(); + + VK_CHECK_RESULT(vkCreateInstance(&create_info, nullptr, &_instance)); +} + +void VulkanTest::CreateDevice() { + UNSCOPED_INFO("Not all of the required device extensions are supported"); + REQUIRE(CheckExtensionSupport(_required_device_extensions)); + + FindPhysicalDevice(); + + VkDeviceQueueCreateInfo queue_create_info = {}; + queue_create_info.sType = VK_STRUCTURE_TYPE_DEVICE_QUEUE_CREATE_INFO; + queue_create_info.queueFamilyIndex = _compute_family_queue_idx = GetComputeQueueFamilyIndex(); + queue_create_info.queueCount = 1; + float queue_priorities = 1.0; + queue_create_info.pQueuePriorities = &queue_priorities; + + VkPhysicalDeviceVulkan12Features features = {}; + features.sType = VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_VULKAN_1_2_FEATURES; + features.timelineSemaphore = true; + + VkDeviceCreateInfo device_create_info = {}; + device_create_info.sType = VK_STRUCTURE_TYPE_DEVICE_CREATE_INFO; + device_create_info.enabledLayerCount = _enabled_layers.size(); + device_create_info.ppEnabledLayerNames = _enabled_layers.data(); + device_create_info.enabledExtensionCount = _required_device_extensions.size(); + device_create_info.ppEnabledExtensionNames = _required_device_extensions.data(); + device_create_info.pQueueCreateInfos = &queue_create_info; + device_create_info.queueCreateInfoCount = 1; + device_create_info.pNext = &features; + + VK_CHECK_RESULT(vkCreateDevice(_physical_device, &device_create_info, nullptr, &_device)); + vkGetDeviceQueue(_device, _compute_family_queue_idx, 0, &_queue); +} + +void VulkanTest::CreateCommandBuffer() { + VkCommandPoolCreateInfo command_pool_create_info = {}; + command_pool_create_info.sType = VK_STRUCTURE_TYPE_COMMAND_POOL_CREATE_INFO; + command_pool_create_info.flags = 0; + command_pool_create_info.queueFamilyIndex = _compute_family_queue_idx; + VK_CHECK_RESULT(vkCreateCommandPool(_device, &command_pool_create_info, nullptr, &_command_pool)); + + VkCommandBufferAllocateInfo command_buffer_allocate_info = {}; + command_buffer_allocate_info.sType = VK_STRUCTURE_TYPE_COMMAND_BUFFER_ALLOCATE_INFO; + command_buffer_allocate_info.commandPool = _command_pool; + command_buffer_allocate_info.level = VK_COMMAND_BUFFER_LEVEL_PRIMARY; + command_buffer_allocate_info.commandBufferCount = 1; + VK_CHECK_RESULT( + vkAllocateCommandBuffers(_device, &command_buffer_allocate_info, &_command_buffer)); +} + +bool VulkanTest::CheckExtensionSupport(std::vector expected_extensions) { + uint32_t extension_count = 0; + vkEnumerateInstanceExtensionProperties(nullptr, &extension_count, nullptr); + std::vector extension_properties(extension_count); + vkEnumerateInstanceExtensionProperties(nullptr, &extension_count, extension_properties.data()); + + std::vector supported_extensions; + supported_extensions.reserve(extension_count); + std::transform(extension_properties.begin(), extension_properties.end(), + std::back_inserter(supported_extensions), + [](const auto& p) { return p.extensionName; }); + + constexpr auto p = [](const char* l, const char* r) { return strcmp(l, r) < 0; }; + std::sort(expected_extensions.begin(), expected_extensions.end(), p); + std::sort(supported_extensions.begin(), supported_extensions.end(), p); + + return std::includes(supported_extensions.begin(), supported_extensions.end(), + expected_extensions.begin(), expected_extensions.end(), + [](const char* l, const char* r) { return strcmp(l, r) == 0; }); +} + +void VulkanTest::EnableValidationLayer() { + uint32_t layer_count = 0; + vkEnumerateInstanceLayerProperties(&layer_count, nullptr); + std::vector layer_properties(layer_count); + vkEnumerateInstanceLayerProperties(&layer_count, layer_properties.data()); + const bool found_val_layer = + std::any_of(layer_properties.cbegin(), layer_properties.cend(), [](const auto& props) { + return strcmp(props.layerName, "VK_LAYER_KHRONOS_validation") == 0; + }); + + + if (found_val_layer) { + _enabled_layers.push_back("VK_LAYER_KHRONOS_validation"); + } else { + UNSCOPED_INFO("Validation was requested, but the validation layer could not be located"); + REQUIRE(found_val_layer); + } +} + +uint32_t VulkanTest::GetComputeQueueFamilyIndex() { + uint32_t queue_family_count = 0u; + + vkGetPhysicalDeviceQueueFamilyProperties(_physical_device, &queue_family_count, nullptr); + std::vector queue_families(queue_family_count); + vkGetPhysicalDeviceQueueFamilyProperties(_physical_device, &queue_family_count, + queue_families.data()); + + const auto it = + std::find_if(queue_families.cbegin(), queue_families.cend(), [](const auto& props) { + return props.queueCount > 0 && (props.queueFlags & VK_QUEUE_COMPUTE_BIT); + }); + REQUIRE(it != queue_families.cend()); + + return std::distance(queue_families.cbegin(), it); +} + +void VulkanTest::FindPhysicalDevice() { + uint32_t device_count = 0; + vkEnumeratePhysicalDevices(_instance, &device_count, nullptr); + REQUIRE(device_count != 0u); + + std::vector physical_devices(device_count); + vkEnumeratePhysicalDevices(_instance, &device_count, physical_devices.data()); + + _physical_device = physical_devices[0]; +} + +uint32_t VulkanTest::FindMemoryType(uint32_t memory_type_bits, VkMemoryPropertyFlags properties) { + VkPhysicalDeviceMemoryProperties memory_properties; + vkGetPhysicalDeviceMemoryProperties(_physical_device, &memory_properties); + for (uint32_t i = 0; i < memory_properties.memoryTypeCount; ++i) { + if ((memory_type_bits & (1 << i)) && + ((memory_properties.memoryTypes[i].propertyFlags & properties) == properties)) { + return i; + } + } + return VK_MAX_MEMORY_TYPES; +} + +hipExternalSemaphoreHandleType VulkanTest::VulkanSemHandleTypeToHIPHandleType( + VkSemaphoreType sem_type) { + if (sem_type == VK_SEMAPHORE_TYPE_BINARY) { + if (_sem_handle_type & VK_EXTERNAL_SEMAPHORE_HANDLE_TYPE_OPAQUE_WIN32_BIT) { + return hipExternalSemaphoreHandleTypeOpaqueWin32; + } else if (_sem_handle_type & VK_EXTERNAL_SEMAPHORE_HANDLE_TYPE_OPAQUE_WIN32_KMT_BIT) { + return hipExternalSemaphoreHandleTypeOpaqueWin32Kmt; + } else if (_sem_handle_type & VK_EXTERNAL_SEMAPHORE_HANDLE_TYPE_OPAQUE_FD_BIT) { + return hipExternalSemaphoreHandleTypeOpaqueFd; + } + } else if (sem_type == VK_SEMAPHORE_TYPE_TIMELINE) { +#if HT_AMD + throw std::invalid_argument("Timeline semaphore unsupported on AMD"); +#else + if (_sem_handle_type & VK_EXTERNAL_SEMAPHORE_HANDLE_TYPE_OPAQUE_WIN32_BIT) { + return hipExternalSemaphoreHandleTypeTimelineSemaphoreWin32; + } else if (_sem_handle_type & VK_EXTERNAL_SEMAPHORE_HANDLE_TYPE_OPAQUE_WIN32_KMT_BIT) { + return hipExternalSemaphoreHandleTypeTimelineSemaphoreWin32; + } else if (_sem_handle_type & VK_EXTERNAL_SEMAPHORE_HANDLE_TYPE_OPAQUE_FD_BIT) { + return hipExternalSemaphoreHandleTypeTimelineSemaphoreFd; + } +#endif + } + + throw std::invalid_argument("Invalid vulkan semaphore handle type"); +} + +hipExternalMemoryHandleType VulkanTest::VulkanMemHandleTypeToHIPHandleType() { + if (_mem_handle_type & VK_EXTERNAL_MEMORY_HANDLE_TYPE_OPAQUE_WIN32_BIT) { + return hipExternalMemoryHandleTypeOpaqueWin32; + } else if (_mem_handle_type & VK_EXTERNAL_MEMORY_HANDLE_TYPE_OPAQUE_WIN32_KMT_BIT) { + return hipExternalMemoryHandleTypeOpaqueWin32Kmt; + } else if (_mem_handle_type & VK_EXTERNAL_MEMORY_HANDLE_TYPE_OPAQUE_FD_BIT) { + return hipExternalMemoryHandleTypeOpaqueFd; + } + + throw std::invalid_argument("Invalid vulkan memory handle type"); +} + +#ifdef _WIN64 +HANDLE +VulkanTest::GetSemaphoreHandle(VkSemaphore semaphore) { + HANDLE handle = 0; + + VkSemaphoreGetWin32HandleInfoKHR semaphoreGetWin32HandleInfoKHR = {}; + semaphoreGetWin32HandleInfoKHR.sType = VK_STRUCTURE_TYPE_SEMAPHORE_GET_WIN32_HANDLE_INFO_KHR; + semaphoreGetWin32HandleInfoKHR.pNext = NULL; + semaphoreGetWin32HandleInfoKHR.semaphore = semaphore; + semaphoreGetWin32HandleInfoKHR.handleType = _sem_handle_type; + + PFN_vkGetSemaphoreWin32HandleKHR fpGetSemaphoreWin32HandleKHR; + fpGetSemaphoreWin32HandleKHR = (PFN_vkGetSemaphoreWin32HandleKHR)vkGetDeviceProcAddr( + _device, "vkGetSemaphoreWin32HandleKHR"); + if (!fpGetSemaphoreWin32HandleKHR) { + throw std::runtime_error("Failed to retrieve vkGetSemaphoreWin32HandleKHR"); + } + if (fpGetSemaphoreWin32HandleKHR(_device, &semaphoreGetWin32HandleInfoKHR, &handle) != + VK_SUCCESS) { + throw std::runtime_error("Failed to retrieve handle for buffer!"); + } + + return handle; +} +#else +int VulkanTest::GetSemaphoreHandle(VkSemaphore semaphore) { + int fd; + + VkSemaphoreGetFdInfoKHR semaphoreGetFdInfoKHR = {}; + semaphoreGetFdInfoKHR.sType = VK_STRUCTURE_TYPE_SEMAPHORE_GET_FD_INFO_KHR; + semaphoreGetFdInfoKHR.pNext = NULL; + semaphoreGetFdInfoKHR.semaphore = semaphore; + semaphoreGetFdInfoKHR.handleType = _sem_handle_type; + + PFN_vkGetSemaphoreFdKHR fpGetSemaphoreFdKHR; + fpGetSemaphoreFdKHR = + (PFN_vkGetSemaphoreFdKHR)vkGetDeviceProcAddr(_device, "vkGetSemaphoreFdKHR"); + if (!fpGetSemaphoreFdKHR) { + throw std::runtime_error("Failed to retrieve vkGetSemaphoreFdKHR"); + } + if (fpGetSemaphoreFdKHR(_device, &semaphoreGetFdInfoKHR, &fd) != VK_SUCCESS) { + throw std::runtime_error("Failed to retrieve semaphore handle"); + } + + return fd; +} +#endif + +#ifdef _WIN64 +HANDLE +VulkanTest::GetMemoryHandle(VkDeviceMemory memory) { + Handle handle = 0; + + VkMemoryGetWin32HandleInfoKHR vkMemoryGetWin32HandleInfoKHR = {}; + vkMemoryGetWin32HandleInfoKHR.sType = VK_STRUCTURE_TYPE_MEMORY_GET_WIN32_HANDLE_INFO_KHR; + vkMemoryGetWin32HandleInfoKHR.memory = memory; + vkMemoryGetWin32HandleInfoKHR.handleType = _mem_handle_type; + + PFN_vkGetMemoryWin32HandleKHR fpGetMemoryWin32HandleKHR = + (PFN_vkGetMemoryWin32HandleKHR)vkGetDeviceProcAddr(m_device, "vkGetMemoryWin32HandleKHR"); + + if (!fpGetMemoryWin32HandleKHR) { + throw std::runtime_error("Failed to retrieve vkGetMemoryWin32HandleKHR"); + } + if (fpGetMemoryWin32HandleKHR(_device, &vkMemoryGetWin32HandleInfoKHR, &handle) != VK_SUCCESS) { + throw std::runtime_error("Failed to retrieve memory handle"); + } + + return handle; +} +#else +int VulkanTest::GetMemoryHandle(VkDeviceMemory memory) { + int fd; + + VkMemoryGetFdInfoKHR memoryGetFdInfoKHR = {}; + memoryGetFdInfoKHR.sType = VK_STRUCTURE_TYPE_MEMORY_GET_FD_INFO_KHR; + memoryGetFdInfoKHR.memory = memory; + memoryGetFdInfoKHR.handleType = _mem_handle_type; + + PFN_vkGetMemoryFdKHR fpGetMemoryFdKHR = + (PFN_vkGetMemoryFdKHR)vkGetDeviceProcAddr(_device, "vkGetMemoryFdKHR"); + if (!fpGetMemoryFdKHR) { + throw std::runtime_error("Failed to retrieve vkGetMemoryFdKHR"); + } + if (fpGetMemoryFdKHR(_device, &memoryGetFdInfoKHR, &fd) != VK_SUCCESS) { + throw std::runtime_error("Failed to retrieve memory handle"); + } + + return fd; +} +#endif + +VkExternalSemaphoreHandleTypeFlagBits VulkanTest::GetVkSemHandlePlatformType() const { +#ifdef _WIN64 + return IsWindows8OrGreater() ? VK_EXTERNAL_SEMAPHORE_HANDLE_TYPE_OPAQUE_WIN32_BIT + : VK_EXTERNAL_SEMAPHORE_HANDLE_TYPE_OPAQUE_WIN32_KMT_BIT; +#else + return VK_EXTERNAL_SEMAPHORE_HANDLE_TYPE_OPAQUE_FD_BIT; +#endif +} + +VkExternalMemoryHandleTypeFlagBits VulkanTest::GetVkMemHandlePlatformType() const { +#ifdef _WIN64 + return IsWindows8OrGreater() ? VK_EXTERNAL_MEMORY_HANDLE_TYPE_OPAQUE_WIN32_BIT + : VK_EXTERNAL_MEMORY_HANDLE_TYPE_OPAQUE_WIN32_KMT_BIT; +#else + return VK_EXTERNAL_MEMORY_HANDLE_TYPE_OPAQUE_FD_BIT; +#endif +} + +// Sometimes in CUDA the stream is not immediately ready after a semaphore has been signaled +void PollStream(hipStream_t stream, hipError_t expected, uint32_t num_iterations) { + hipError_t query_result; + for (uint32_t _ = 0; _ < num_iterations; ++_) { + if ((query_result = hipStreamQuery(stream)) != expected) { + std::this_thread::sleep_for(std::chrono::milliseconds{5}); + } else { + break; + } + } + REQUIRE(expected == query_result); +} + +hipExternalSemaphore_t ImportBinarySemaphore(VulkanTest& vkt) { + const auto semaphore = vkt.CreateExternalSemaphore(VK_SEMAPHORE_TYPE_BINARY); + const auto sem_handle_desc = vkt.BuildSemaphoreDescriptor(semaphore, VK_SEMAPHORE_TYPE_BINARY); + hipExternalSemaphore_t hip_ext_semaphore; + HIP_CHECK(hipImportExternalSemaphore(&hip_ext_semaphore, &sem_handle_desc)); + + return hip_ext_semaphore; +} \ No newline at end of file diff --git a/projects/hip-tests/catch/unit/vulkan_interop/vulkan_test.hh b/projects/hip-tests/catch/unit/vulkan_interop/vulkan_test.hh new file mode 100644 index 0000000000..e223e8cd2b --- /dev/null +++ b/projects/hip-tests/catch/unit/vulkan_interop/vulkan_test.hh @@ -0,0 +1,259 @@ +/* +Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved. +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#pragma once + +#include +#include + +#ifdef _WIN64 +#include +#endif + +#include +#include + +#define VK_CHECK_RESULT(code) \ + { \ + VkResult res = (code); \ + if (res != VK_SUCCESS) { \ + INFO("Vulkan error: " << std::to_string(res) << "\n In File: " << __FILE__ \ + << "\n At line: " << __LINE__); \ + REQUIRE(false); \ + } \ + } + +class VulkanTest { + public: + VulkanTest(bool enable_validation) + : _enable_validation{enable_validation}, + _sem_handle_type{GetVkSemHandlePlatformType()}, + _mem_handle_type{GetVkMemHandlePlatformType()} { + CreateInstance(); + CreateDevice(); + CreateCommandBuffer(); + } + + ~VulkanTest() { + for (const auto s : _semaphores) { + vkDestroySemaphore(_device, s, nullptr); + } + + for (const auto f : _fences) { + vkDestroyFence(_device, f, nullptr); + } + + for (const auto& s : _stores) { + vkUnmapMemory(_device, s.memory); + vkDestroyBuffer(_device, s.buffer, nullptr); + vkFreeMemory(_device, s.memory, nullptr); + } + + if (_command_buffer != VK_NULL_HANDLE) + vkFreeCommandBuffers(_device, _command_pool, 1, &_command_buffer); + + if (_command_pool != VK_NULL_HANDLE) vkDestroyCommandPool(_device, _command_pool, nullptr); + + if (_device != VK_NULL_HANDLE) vkDestroyDevice(_device, nullptr); + + if (_instance != VK_NULL_HANDLE) vkDestroyInstance(_instance, nullptr); + } + + VulkanTest(const VulkanTest&) = delete; + + VulkanTest(VulkanTest&&) = delete; + + template struct MappedBuffer { + VkDeviceMemory memory = VK_NULL_HANDLE; + VkBuffer buffer = VK_NULL_HANDLE; + uint32_t size = 0; + T* host_ptr = nullptr; + }; + + template + MappedBuffer CreateMappedStorage(uint32_t count, VkBufferUsageFlagBits transfer_flags, + bool external = false); + + VkFence CreateFence(); + + VkSemaphore CreateExternalSemaphore(VkSemaphoreType sem_type, uint64_t initial_value = 0); + + hipExternalSemaphoreHandleDesc BuildSemaphoreDescriptor(VkSemaphore vk_sem, + VkSemaphoreType sem_type); + + hipExternalMemoryHandleDesc BuildMemoryDescriptor(VkDeviceMemory vk_mem, uint32_t size); + + + VkDevice GetDevice() const { return _device; } + + VkCommandBuffer GetCommandBuffer() const { return _command_buffer; } + + VkQueue GetQueue() const { return _queue; } + + private: + void CreateInstance(); + + void CreateDevice(); + + void CreateCommandBuffer(); + + bool CheckExtensionSupport(std::vector expected_extensions); + + void EnableValidationLayer(); + + uint32_t GetComputeQueueFamilyIndex(); + + void FindPhysicalDevice(); + + uint32_t FindMemoryType(uint32_t memory_type_bits, VkMemoryPropertyFlags properties); + + hipExternalSemaphoreHandleType VulkanSemHandleTypeToHIPHandleType(VkSemaphoreType sem_type); + + hipExternalMemoryHandleType VulkanMemHandleTypeToHIPHandleType(); + +#ifdef _WIN64 + HANDLE + GetSemaphoreHandle(VkSemaphore semaphore); +#else + int GetSemaphoreHandle(VkSemaphore semaphore); +#endif + +#ifdef _WIN64 + HANDLE + GetMemoryHandle(VkDeviceMemory memory); +#else + int GetMemoryHandle(VkDeviceMemory memory); +#endif + + VkExternalSemaphoreHandleTypeFlagBits GetVkSemHandlePlatformType() const; + + VkExternalMemoryHandleTypeFlagBits GetVkMemHandlePlatformType() const; + + struct Storage { + VkBuffer buffer = VK_NULL_HANDLE; + VkDeviceMemory memory = VK_NULL_HANDLE; + uint32_t size = 0u; + }; + + private: + const bool _enable_validation = false; + const VkExternalSemaphoreHandleTypeFlagBits _sem_handle_type; + const VkExternalMemoryHandleTypeFlagBits _mem_handle_type; + VkInstance _instance = VK_NULL_HANDLE; + VkPhysicalDevice _physical_device = VK_NULL_HANDLE; + VkDevice _device = VK_NULL_HANDLE; + VkQueue _queue = VK_NULL_HANDLE; + VkCommandPool _command_pool = VK_NULL_HANDLE; + VkCommandBuffer _command_buffer = VK_NULL_HANDLE; + uint32_t _compute_family_queue_idx = 0u; + std::vector _enabled_layers; + + std::vector _semaphores; + std::vector _fences; + std::vector _stores; + + std::vector _required_instance_extensions{ + VK_KHR_GET_PHYSICAL_DEVICE_PROPERTIES_2_EXTENSION_NAME, + VK_KHR_EXTERNAL_SEMAPHORE_CAPABILITIES_EXTENSION_NAME, + VK_KHR_EXTERNAL_MEMORY_CAPABILITIES_EXTENSION_NAME}; +#ifdef _WIN64 + std::vector _required_device_extensions{ + VK_KHR_EXTERNAL_SEMAPHORE_EXTENSION_NAME, VK_KHR_EXTERNAL_SEMAPHORE_WIN32_EXTENSION_NAME, + VK_KHR_EXTERNAL_MEMORY_EXTENSION_NAME, VK_KHR_EXTERNAL_MEMORY_WIN32_EXTENSION_NAME}; +#else + std::vector _required_device_extensions{ + VK_KHR_EXTERNAL_SEMAPHORE_EXTENSION_NAME, VK_KHR_EXTERNAL_SEMAPHORE_FD_EXTENSION_NAME, + VK_KHR_EXTERNAL_MEMORY_EXTENSION_NAME, VK_KHR_EXTERNAL_MEMORY_FD_EXTENSION_NAME}; +#endif +}; + + +template +VulkanTest::MappedBuffer VulkanTest::CreateMappedStorage(uint32_t count, + VkBufferUsageFlagBits transfer_flags, + bool external) { + Storage storage; + const auto size = count * sizeof(T); + + VkBufferCreateInfo buffer_create_info = {}; + buffer_create_info.sType = VK_STRUCTURE_TYPE_BUFFER_CREATE_INFO; + buffer_create_info.size = size; + buffer_create_info.usage = VK_BUFFER_USAGE_STORAGE_BUFFER_BIT | transfer_flags; + buffer_create_info.sharingMode = VK_SHARING_MODE_EXCLUSIVE; + + VkExternalMemoryBufferCreateInfo external_memory_buffer_info = {}; + if (external) { + external_memory_buffer_info.sType = VK_STRUCTURE_TYPE_EXTERNAL_MEMORY_BUFFER_CREATE_INFO; + external_memory_buffer_info.handleTypes = _mem_handle_type; + buffer_create_info.pNext = &external_memory_buffer_info; + } + VK_CHECK_RESULT(vkCreateBuffer(_device, &buffer_create_info, nullptr, &storage.buffer)); + + VkMemoryRequirements memory_requirements; + vkGetBufferMemoryRequirements(_device, storage.buffer, &memory_requirements); + storage.size = memory_requirements.size; + + VkMemoryAllocateInfo allocate_info = {}; + allocate_info.sType = VK_STRUCTURE_TYPE_MEMORY_ALLOCATE_INFO; + allocate_info.allocationSize = memory_requirements.size; + allocate_info.memoryTypeIndex = + FindMemoryType(memory_requirements.memoryTypeBits, + VK_MEMORY_PROPERTY_HOST_COHERENT_BIT | VK_MEMORY_PROPERTY_HOST_VISIBLE_BIT); + REQUIRE(allocate_info.memoryTypeIndex != VK_MAX_MEMORY_TYPES); + + VkExportMemoryAllocateInfoKHR vulkan_export_memory_allocate_info = {}; + if (external) { + vulkan_export_memory_allocate_info.sType = VK_STRUCTURE_TYPE_EXPORT_MEMORY_ALLOCATE_INFO_KHR; + vulkan_export_memory_allocate_info.handleTypes = _mem_handle_type; + +#ifdef _WIN64 + WindowsSecurityAttributes winSecurityAttributes; + + VkExportMemoryWin32HandleInfoKHR vulkanExportMemoryWin32HandleInfoKHR = {}; + vulkanExportMemoryWin32HandleInfoKHR.sType = + VK_STRUCTURE_TYPE_EXPORT_MEMORY_WIN32_HANDLE_INFO_KHR; + vulkanExportMemoryWin32HandleInfoKHR.pNext = NULL; + vulkanExportMemoryWin32HandleInfoKHR.pAttributes = &winSecurityAttributes; + vulkanExportMemoryWin32HandleInfoKHR.dwAccess = + DXGI_SHARED_RESOURCE_READ | DXGI_SHARED_RESOURCE_WRITE; + vulkanExportMemoryWin32HandleInfoKHR.name = (LPCWSTR)NULL; + + vulkan_export_memory_allocate_info.pNext = + _mem_handle_type & VK_EXTERNAL_MEMORY_HANDLE_TYPE_OPAQUE_WIN32_BIT_KHR + ? &vulkanExportMemoryWin32HandleInfoKHR + : NULL; +#endif + allocate_info.pNext = &vulkan_export_memory_allocate_info; + } + + VK_CHECK_RESULT(vkAllocateMemory(_device, &allocate_info, nullptr, &storage.memory)); + VK_CHECK_RESULT(vkBindBufferMemory(_device, storage.buffer, storage.memory, 0)); + + T* host_ptr = nullptr; + VK_CHECK_RESULT(vkMapMemory(_device, storage.memory, 0, storage.size, 0, + reinterpret_cast(&host_ptr))); + + _stores.push_back(storage); + return MappedBuffer{storage.memory, storage.buffer, storage.size, host_ptr}; +} + +// Sometimes in CUDA the stream is not immediately ready after a semaphore has been signaled +void PollStream(hipStream_t stream, hipError_t expected, uint32_t num_iterations = 5); + +hipExternalSemaphore_t ImportBinarySemaphore(VulkanTest& vkt); \ No newline at end of file