EXSWHTEC-149 - Implement tests for hipExternalMemoryGetMappedBuffer for the Vulkan API (#25)
- Basic positive test
- Negative parameter tests
[ROCm/hip-tests commit: a3e1dcd26c]
This commit is contained in:
music-dino
GitHub
@@ -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()
|
||||
|
||||
@@ -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)
|
||||
@@ -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 <typename T> __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<type>(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<void**>(&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<type> 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<type>(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<type>(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<void**>(&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<int>(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
|
||||
}
|
||||
@@ -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 <iostream>
|
||||
#include <algorithm>
|
||||
|
||||
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<uint32_t>(_required_instance_extensions.size());
|
||||
create_info.ppEnabledExtensionNames = _required_instance_extensions.data();
|
||||
create_info.enabledLayerCount = static_cast<uint32_t>(_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<const char*> expected_extensions) {
|
||||
uint32_t extension_count = 0;
|
||||
vkEnumerateInstanceExtensionProperties(nullptr, &extension_count, nullptr);
|
||||
std::vector<VkExtensionProperties> extension_properties(extension_count);
|
||||
vkEnumerateInstanceExtensionProperties(nullptr, &extension_count, extension_properties.data());
|
||||
|
||||
std::vector<const char*> 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<VkLayerProperties> 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<VkQueueFamilyProperties> 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<VkPhysicalDevice> 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;
|
||||
}
|
||||
@@ -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 <vulkan/vulkan.h>
|
||||
#include <vector>
|
||||
|
||||
#ifdef _WIN64
|
||||
#include <VersionHelpers.h>
|
||||
#endif
|
||||
|
||||
#include <hip_test_common.hh>
|
||||
#include <hip/hip_runtime_api.h>
|
||||
|
||||
#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 <typename T> struct MappedBuffer {
|
||||
VkDeviceMemory memory = VK_NULL_HANDLE;
|
||||
VkBuffer buffer = VK_NULL_HANDLE;
|
||||
uint32_t size = 0;
|
||||
T* host_ptr = nullptr;
|
||||
};
|
||||
|
||||
template <typename T>
|
||||
MappedBuffer<T> 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<const char*> 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<const char*> _enabled_layers;
|
||||
|
||||
std::vector<VkSemaphore> _semaphores;
|
||||
std::vector<VkFence> _fences;
|
||||
std::vector<Storage> _stores;
|
||||
|
||||
std::vector<const char*> _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<const char*> _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<const char*> _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 <typename T>
|
||||
VulkanTest::MappedBuffer<T> 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<void**>(&host_ptr)));
|
||||
|
||||
_stores.push_back(storage);
|
||||
return MappedBuffer<T>{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);
|
||||
Reference in New Issue
Block a user