SWDEV-558855 - hipExternalMemoryGetMappedBuffer test with CPU-nonvisible memory (#1760)
Este commit está contenido en:
Jin Jung
GitHub
+130
-1
@@ -194,4 +194,133 @@ TEST_CASE("Unit_hipExternalMemoryGetMappedBuffer_Vulkan_Capture") {
|
||||
memcpy_err);
|
||||
END_CAPTURE_SYNC(memcpy_err);
|
||||
REQUIRE(nullptr != hip_dev_ptr);
|
||||
}
|
||||
}
|
||||
|
||||
TEST_CASE("Unit_hipExternalMemoryGetMappedBuffer_Vulkan_Positive_Read_Write_Device_Memory") {
|
||||
VulkanTest vkt(enable_validation);
|
||||
using type = uint8_t;
|
||||
constexpr uint32_t count = 3;
|
||||
|
||||
uint32_t size = count * sizeof(type);
|
||||
VkDeviceMemory memory = VK_NULL_HANDLE;
|
||||
VkBuffer buffer = VK_NULL_HANDLE;
|
||||
|
||||
vkt.CreateBuffer(size, VK_BUFFER_USAGE_TRANSFER_DST_BIT | VK_BUFFER_USAGE_STORAGE_BUFFER_BIT,
|
||||
VK_MEMORY_PROPERTY_DEVICE_LOCAL_BIT,
|
||||
buffer, memory, true);
|
||||
|
||||
if (memory == nullptr) {
|
||||
return;
|
||||
}
|
||||
const auto hip_ext_mem_desc = vkt.BuildMemoryDescriptor(memory, size);
|
||||
|
||||
// Staging buffer creation and data copy to/from Vulkan buffer.
|
||||
VkBuffer src_staging_buffer = VK_NULL_HANDLE;
|
||||
VkDeviceMemory src_staging_memory = VK_NULL_HANDLE;
|
||||
type* src_data;
|
||||
vkt.CreateBuffer(size, VK_BUFFER_USAGE_TRANSFER_SRC_BIT,
|
||||
VK_MEMORY_PROPERTY_HOST_VISIBLE_BIT | VK_MEMORY_PROPERTY_HOST_COHERENT_BIT,
|
||||
src_staging_buffer, src_staging_memory);
|
||||
vkMapMemory(vkt.GetDevice(), src_staging_memory, 0, size, 0,
|
||||
reinterpret_cast<void**>(&src_data));
|
||||
|
||||
VkBuffer dst_staging_buffer = VK_NULL_HANDLE;
|
||||
VkDeviceMemory dst_staging_memory = VK_NULL_HANDLE;
|
||||
type* dst_data;
|
||||
vkt.CreateBuffer(size, VK_BUFFER_USAGE_TRANSFER_DST_BIT,
|
||||
VK_MEMORY_PROPERTY_HOST_VISIBLE_BIT | VK_MEMORY_PROPERTY_HOST_COHERENT_BIT,
|
||||
dst_staging_buffer, dst_staging_memory);
|
||||
vkMapMemory(vkt.GetDevice(), dst_staging_memory, 0, size, 0,
|
||||
reinterpret_cast<void**>(&dst_data));
|
||||
|
||||
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 = 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);
|
||||
|
||||
src_data[0] = 41;
|
||||
src_data[1] = 40;
|
||||
src_data[2] = 43;
|
||||
|
||||
vkt.CopyBuffer(src_staging_buffer, buffer, size);
|
||||
|
||||
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());
|
||||
|
||||
vkt.CopyBuffer(buffer, dst_staging_buffer, size);
|
||||
|
||||
REQUIRE(41 == dst_data[0]);
|
||||
REQUIRE(42 == dst_data[1]);
|
||||
REQUIRE(43 == dst_data[2]);
|
||||
|
||||
HIP_CHECK(hipFree(hip_dev_ptr));
|
||||
HIP_CHECK(hipDestroyExternalMemory(hip_ext_memory));
|
||||
|
||||
vkDestroyBuffer(vkt.GetDevice(), buffer, nullptr);
|
||||
vkFreeMemory(vkt.GetDevice(), memory, nullptr);
|
||||
|
||||
vkUnmapMemory(vkt.GetDevice(), src_staging_memory);
|
||||
vkDestroyBuffer(vkt.GetDevice(), src_staging_buffer, nullptr);
|
||||
vkFreeMemory(vkt.GetDevice(), src_staging_memory, nullptr);
|
||||
|
||||
vkUnmapMemory(vkt.GetDevice(), dst_staging_memory);
|
||||
vkDestroyBuffer(vkt.GetDevice(), dst_staging_buffer, nullptr);
|
||||
vkFreeMemory(vkt.GetDevice(), dst_staging_memory, nullptr);
|
||||
}
|
||||
|
||||
TEST_CASE("Unit_hipExternalMemoryGetMappedBuffer_Vulkan_Positive_Read_Write_With_Offset_Device_Memory") {
|
||||
VulkanTest vkt(enable_validation);
|
||||
using type = uint8_t;
|
||||
constexpr uint32_t count = 2;
|
||||
|
||||
uint32_t size = count * sizeof(type);
|
||||
VkDeviceMemory memory = VK_NULL_HANDLE;
|
||||
VkBuffer buffer = VK_NULL_HANDLE;
|
||||
|
||||
vkt.CreateBuffer(size, VK_BUFFER_USAGE_TRANSFER_DST_BIT | VK_BUFFER_USAGE_STORAGE_BUFFER_BIT,
|
||||
VK_MEMORY_PROPERTY_DEVICE_LOCAL_BIT,
|
||||
buffer, memory, true);
|
||||
|
||||
if (memory == nullptr) {
|
||||
return;
|
||||
}
|
||||
const auto hip_ext_mem_desc = vkt.BuildMemoryDescriptor(memory, 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 = 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));
|
||||
|
||||
Set<<<1, 1>>>(hip_dev_ptr, static_cast<type>(42));
|
||||
HIP_CHECK(hipDeviceSynchronize());
|
||||
|
||||
type read_val = 0;
|
||||
HIP_CHECK(hipMemcpy(&read_val, hip_dev_ptr, 1, hipMemcpyDeviceToHost));
|
||||
REQUIRE(42 == read_val);
|
||||
|
||||
HIP_CHECK(hipFree(hip_dev_ptr));
|
||||
HIP_CHECK(hipDestroyExternalMemory(hip_ext_memory));
|
||||
|
||||
vkDestroyBuffer(vkt.GetDevice(), buffer, nullptr);
|
||||
vkFreeMemory(vkt.GetDevice(), memory, nullptr);
|
||||
}
|
||||
|
||||
@@ -389,6 +389,96 @@ VkExternalMemoryHandleTypeFlagBits VulkanTest::GetVkMemHandlePlatformType() cons
|
||||
#endif
|
||||
}
|
||||
|
||||
void VulkanTest::CreateBuffer(VkDeviceSize size, VkBufferUsageFlags usage,
|
||||
VkMemoryPropertyFlags properties, VkBuffer& buffer,
|
||||
VkDeviceMemory& buffer_memory, bool external) {
|
||||
VkBufferCreateInfo buffer_create_info = {};
|
||||
buffer_create_info.sType = VK_STRUCTURE_TYPE_BUFFER_CREATE_INFO;
|
||||
buffer_create_info.size = size;
|
||||
buffer_create_info.usage = usage;
|
||||
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, &buffer));
|
||||
|
||||
VkMemoryRequirements memory_requirements;
|
||||
vkGetBufferMemoryRequirements(_device, buffer, &memory_requirements);
|
||||
|
||||
VkMemoryAllocateInfo alloc_info = {};
|
||||
alloc_info.sType = VK_STRUCTURE_TYPE_MEMORY_ALLOCATE_INFO;
|
||||
alloc_info.allocationSize = memory_requirements.size;
|
||||
alloc_info.memoryTypeIndex = FindMemoryType(memory_requirements.memoryTypeBits, properties);
|
||||
|
||||
VkExportMemoryAllocateInfoKHR vulkan_export_memory_allocate_info = {};
|
||||
#ifdef _WIN64
|
||||
WindowsSecurityAttributes win_security_attributes = {};
|
||||
VkExportMemoryWin32HandleInfoKHR vulkan_export_memory_win32_handle_info = {};
|
||||
#endif
|
||||
|
||||
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
|
||||
vulkan_export_memory_win32_handle_info.sType =
|
||||
VK_STRUCTURE_TYPE_EXPORT_MEMORY_WIN32_HANDLE_INFO_KHR;
|
||||
vulkan_export_memory_win32_handle_info.pNext = NULL;
|
||||
vulkan_export_memory_win32_handle_info.pAttributes = &win_security_attributes;
|
||||
vulkan_export_memory_win32_handle_info.dwAccess =
|
||||
DXGI_SHARED_RESOURCE_READ | DXGI_SHARED_RESOURCE_WRITE;
|
||||
vulkan_export_memory_win32_handle_info.name = (LPCWSTR)NULL;
|
||||
|
||||
vulkan_export_memory_allocate_info.pNext =
|
||||
_mem_handle_type & VK_EXTERNAL_MEMORY_HANDLE_TYPE_OPAQUE_WIN32_BIT_KHR
|
||||
? &vulkan_export_memory_win32_handle_info
|
||||
: NULL;
|
||||
#endif
|
||||
|
||||
alloc_info.pNext = &vulkan_export_memory_allocate_info;
|
||||
}
|
||||
|
||||
VK_CHECK_RESULT(vkAllocateMemory(_device, &alloc_info, nullptr, &buffer_memory));
|
||||
VK_CHECK_RESULT(vkBindBufferMemory(_device, buffer, buffer_memory, 0));
|
||||
}
|
||||
|
||||
void VulkanTest::CopyBuffer(VkBuffer src_buffer, VkBuffer dst_buffer, VkDeviceSize size) {
|
||||
VkCommandBufferAllocateInfo alloc_info = {};
|
||||
alloc_info.sType = VK_STRUCTURE_TYPE_COMMAND_BUFFER_ALLOCATE_INFO;
|
||||
alloc_info.level = VK_COMMAND_BUFFER_LEVEL_PRIMARY;
|
||||
alloc_info.commandPool = _command_pool;
|
||||
alloc_info.commandBufferCount = 1;
|
||||
|
||||
VkCommandBuffer command_buffer;
|
||||
vkAllocateCommandBuffers(_device, &alloc_info, &command_buffer);
|
||||
|
||||
VkCommandBufferBeginInfo begin_info = {};
|
||||
begin_info.sType = VK_STRUCTURE_TYPE_COMMAND_BUFFER_BEGIN_INFO;
|
||||
begin_info.flags = VK_COMMAND_BUFFER_USAGE_ONE_TIME_SUBMIT_BIT;
|
||||
|
||||
vkBeginCommandBuffer(command_buffer, &begin_info);
|
||||
|
||||
VkBufferCopy buffer_copy = {};
|
||||
buffer_copy.size = size;
|
||||
vkCmdCopyBuffer(command_buffer, src_buffer, dst_buffer, 1, &buffer_copy);
|
||||
|
||||
vkEndCommandBuffer(command_buffer);
|
||||
|
||||
VkSubmitInfo submit_info = {};
|
||||
submit_info.sType = VK_STRUCTURE_TYPE_SUBMIT_INFO;
|
||||
submit_info.commandBufferCount = 1;
|
||||
submit_info.pCommandBuffers = &command_buffer;
|
||||
|
||||
vkQueueSubmit(_queue, 1, &submit_info, VK_NULL_HANDLE);
|
||||
vkQueueWaitIdle(_queue);
|
||||
vkFreeCommandBuffers(_device, _command_pool, 1, &command_buffer);
|
||||
}
|
||||
|
||||
// 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;
|
||||
|
||||
@@ -182,6 +182,11 @@ class VulkanTest {
|
||||
|
||||
VkQueue GetQueue() const { return _queue; }
|
||||
|
||||
void CreateBuffer(VkDeviceSize size, VkBufferUsageFlags usage, VkMemoryPropertyFlags properties,
|
||||
VkBuffer& buffer, VkDeviceMemory& buffer_memory, bool external = false);
|
||||
|
||||
void CopyBuffer(VkBuffer src_buffer, VkBuffer dst_buffer, VkDeviceSize size);
|
||||
|
||||
private:
|
||||
void CreateInstance();
|
||||
|
||||
|
||||
Referencia en una nueva incidencia
Block a user