SWDEV-408843 - Fix vulkan_interop failure
Fix vulkan_interop build issue on CMake so that it can be integrated into catch2 batch build. Fix some test failures. Update some code format. Change-Id: I32c9bed861ddf4fe0d7bba21dce9bd720168c397
This commit is contained in:
taosang2
Rakesh Roy
@@ -10,25 +10,41 @@ set(TEST_SRC
|
||||
hipDestroyExternalSemaphore.cc
|
||||
)
|
||||
|
||||
if(WIN32)
|
||||
set(Vulkan_LIBRARY $ENV{VULKAN_SDK}/Lib/vulkan-l)
|
||||
else()
|
||||
# The code can be compiled in Linux, but some required instance extentions
|
||||
# aren't supported in Linux, thus test will fail at VulkanTest::CreateInstance()
|
||||
# with VK_ERROR_EXTENSION_NOT_PRESENT(-7).
|
||||
# Now temperally ignore it in Linux
|
||||
message(STATUS "Ignore vulkan_interop test in Linux")
|
||||
return()
|
||||
endif()
|
||||
find_package(Vulkan)
|
||||
message(STATUS "Vulkan_FOUND: ${Vulkan_FOUND}")
|
||||
message(STATUS "Vulkan_LIBRARIES: ${Vulkan_LIBRARIES}")
|
||||
message(STATUS "Vulkan_INCLUDE_DIRS: ${Vulkan_INCLUDE_DIRS}")
|
||||
|
||||
if(NOT Vulkan_FOUND)
|
||||
if(EXISTS "${VULKAN_PATH}")
|
||||
message(STATUS "Vulkan SDK: ${VULKAN_PATH}")
|
||||
elseif (EXISTS "$ENV{VULKAN_SDK}")
|
||||
if (EXISTS "$ENV{VULKAN_SDK}")
|
||||
message(STATUS "FOUND VULKAN SDK: $ENV{VULKAN_SDK}")
|
||||
set(VULKAN_PATH $ENV{VULKAN_SDK})
|
||||
message(STATUS "Please check Vulkan_LIBRARY if it is missing. Ignore vulkan test!")
|
||||
return()
|
||||
else()
|
||||
message(STATUS "Vulkan sdk not found, interop test not enabled. Please add VULKAN_PATH to enable the test")
|
||||
message(STATUS "Vulkan sdk not found, interop test not enabled. Please set ENV VULKAN_SDK to enable the test")
|
||||
return()
|
||||
endif()
|
||||
endif()
|
||||
|
||||
if(WIN32)
|
||||
set(LINKER_LIBS vulkan-1)
|
||||
set(LINKER_LIBS vulkan-1 advapi32)
|
||||
else()
|
||||
set(LINKER_LIBS vulkan)
|
||||
endif()
|
||||
|
||||
#set(CMAKE_BUILD_TYPE DEBUG)
|
||||
include_directories(AFTER ${Vulkan_INCLUDE_DIRS})
|
||||
link_directories(${Vulkan_INCLUDE_DIRS}/../bin;${Vulkan_INCLUDE_DIRS}/../lib)
|
||||
hip_add_exe_to_target(NAME VulkanInteropTest
|
||||
TEST_SRC ${TEST_SRC}
|
||||
TEST_TARGET_NAME build_tests
|
||||
|
||||
@@ -28,8 +28,8 @@ TEST_CASE("Unit_hipDestroyExternalSemaphore_Vulkan_Negative_Parameters") {
|
||||
HIP_CHECK_ERROR(hipDestroyExternalSemaphore(nullptr), hipErrorInvalidValue);
|
||||
}
|
||||
|
||||
// Segfaults in CUDA
|
||||
#if HT_AMD
|
||||
// Segfaults in Nvidia and Amd
|
||||
#if 0
|
||||
SECTION("Double free") {
|
||||
VulkanTest vkt(enable_validation);
|
||||
const auto ext_semaphore = ImportBinarySemaphore(vkt);
|
||||
|
||||
@@ -32,7 +32,9 @@ TEST_CASE("Unit_hipExternalMemoryGetMappedBuffer_Vulkan_Positive_Read_Write") {
|
||||
|
||||
const auto vk_storage =
|
||||
vkt.CreateMappedStorage<type>(count, VK_BUFFER_USAGE_TRANSFER_DST_BIT, true);
|
||||
|
||||
if (vk_storage.memory == nullptr) {
|
||||
return;
|
||||
}
|
||||
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));
|
||||
@@ -62,13 +64,11 @@ TEST_CASE("Unit_hipExternalMemoryGetMappedBuffer_Vulkan_Positive_Read_Write") {
|
||||
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(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;
|
||||
@@ -76,6 +76,9 @@ TEST_CASE("Unit_hipExternalMemoryGetMappedBuffer_Vulkan_Positive_Read_Write_With
|
||||
|
||||
const auto vk_storage =
|
||||
vkt.CreateMappedStorage<type>(count, VK_BUFFER_USAGE_TRANSFER_DST_BIT, true);
|
||||
if (vk_storage.memory == nullptr) {
|
||||
return;
|
||||
}
|
||||
|
||||
const auto hip_ext_mem_desc = vkt.BuildMemoryDescriptor(vk_storage.memory, vk_storage.size);
|
||||
hipExternalMemory_t hip_ext_memory;
|
||||
@@ -97,15 +100,16 @@ TEST_CASE("Unit_hipExternalMemoryGetMappedBuffer_Vulkan_Positive_Read_Write_With
|
||||
REQUIRE(42 == read_val);
|
||||
|
||||
// Defect - EXSWHTEC-181
|
||||
// HIP_CHECK(hipFree(hip_dev_ptr));
|
||||
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);
|
||||
|
||||
if (vk_storage.memory == nullptr) {
|
||||
return;
|
||||
}
|
||||
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));
|
||||
|
||||
@@ -71,6 +71,8 @@ TEST_CASE("Unit_hipImportExternalMemory_Vulkan_Negative_Parameters") {
|
||||
|
||||
#ifdef _WIN32
|
||||
SECTION("memHandleDesc.handle == NULL") {
|
||||
hipExternalMemory_t ext_memory;
|
||||
hipExternalMemoryHandleDesc desc;
|
||||
desc.handle.win32.handle = NULL;
|
||||
HIP_CHECK_ERROR(hipImportExternalMemory(&ext_memory, &desc), hipErrorInvalidValue);
|
||||
}
|
||||
|
||||
@@ -36,11 +36,13 @@ TEST_CASE("Unit_hipImportExternalSemaphore_Vulkan_Negative_Parameters") {
|
||||
SECTION("semHandleDesc == nullptr") {
|
||||
HIP_CHECK_ERROR(hipImportExternalSemaphore(&ext_semaphore, nullptr), hipErrorInvalidValue);
|
||||
}
|
||||
|
||||
/*
|
||||
* CUDA doesn't specify the case
|
||||
SECTION("semHandleDesc.flags != 0") {
|
||||
handle_desc.flags = 1;
|
||||
HIP_CHECK_ERROR(hipImportExternalSemaphore(&ext_semaphore, &handle_desc), hipErrorInvalidValue);
|
||||
}
|
||||
*/
|
||||
|
||||
SECTION("Invalid semHandleDesc.type") {
|
||||
handle_desc.type = static_cast<hipExternalSemaphoreHandleType>(-1);
|
||||
|
||||
@@ -39,32 +39,31 @@ TEST_CASE("Unit_hipSignalExternalSemaphoresAsync_Vulkan_Positive_Binary_Semaphor
|
||||
buffer_copy.size = count * sizeof(*src_storage.host_ptr);
|
||||
vkCmdCopyBuffer(command_buffer, src_storage.buffer, dst_storage.buffer, 1, &buffer_copy);
|
||||
VK_CHECK_RESULT(vkEndCommandBuffer(command_buffer));
|
||||
|
||||
const auto semaphore = vkt.CreateExternalSemaphore(VK_SEMAPHORE_TYPE_BINARY);
|
||||
const auto hip_sem_handle_desc =
|
||||
vkt.BuildSemaphoreDescriptor(semaphore, VK_SEMAPHORE_TYPE_BINARY);
|
||||
hipExternalSemaphore_t hip_ext_semaphore;
|
||||
HIP_CHECK(hipImportExternalSemaphore(&hip_ext_semaphore, &hip_sem_handle_desc));
|
||||
|
||||
hipExternalSemaphoreSignalParams signal_params = {};
|
||||
signal_params.params.fence.value = 0;
|
||||
HIP_CHECK(hipSignalExternalSemaphoresAsync(&hip_ext_semaphore, &signal_params, 1, nullptr));
|
||||
HIP_CHECK(hipDeviceSynchronize());
|
||||
VkSubmitInfo submit_info = {};
|
||||
submit_info.sType = VK_STRUCTURE_TYPE_SUBMIT_INFO;
|
||||
submit_info.commandBufferCount = 1;
|
||||
submit_info.pCommandBuffers = &command_buffer;
|
||||
VkSemaphore waitSemaphores[] = {semaphore};
|
||||
// VK_PIPELINE_STAGE_TOP_OF_PIPE_BIT or VK_PIPELINE_STAGE_TRANSFER_BIT can work
|
||||
VkPipelineStageFlags waitStages[] = {VK_PIPELINE_STAGE_TOP_OF_PIPE_BIT};
|
||||
submit_info.waitSemaphoreCount = 1;
|
||||
submit_info.pWaitSemaphores = &semaphore;
|
||||
submit_info.pWaitSemaphores = waitSemaphores;
|
||||
submit_info.pWaitDstStageMask = waitStages;
|
||||
const auto fence = vkt.CreateFence();
|
||||
VK_CHECK_RESULT(vkQueueSubmit(vkt.GetQueue(), 1, &submit_info, fence));
|
||||
|
||||
REQUIRE(vkGetFenceStatus(vkt.GetDevice(), fence) == VK_NOT_READY);
|
||||
|
||||
hipExternalSemaphoreSignalParams signal_params = {};
|
||||
signal_params.params.fence.value = 0;
|
||||
HIP_CHECK(hipSignalExternalSemaphoresAsync(&hip_ext_semaphore, &signal_params, 1, nullptr));
|
||||
PollStream(nullptr, hipSuccess);
|
||||
|
||||
VK_CHECK_RESULT(
|
||||
vkWaitForFences(vkt.GetDevice(), 1, &fence, VK_TRUE, 5'000'000'000 /*5 seconds*/));
|
||||
|
||||
HIP_CHECK(hipDestroyExternalSemaphore(hip_ext_semaphore));
|
||||
}
|
||||
|
||||
@@ -93,14 +92,15 @@ TEST_CASE("Unit_hipSignalExternalSemaphoresAsync_Vulkan_Positive_Timeline_Semaph
|
||||
|
||||
HIP_CHECK(hipDestroyExternalSemaphore(hip_ext_semaphore));
|
||||
}
|
||||
#endif
|
||||
|
||||
TEST_CASE("Unit_hipSignalExternalSemaphoresAsync_Vulkan_Positive_Multiple_Semaphores") {
|
||||
VulkanTest vkt(enable_validation);
|
||||
|
||||
constexpr uint32_t count = 1;
|
||||
const auto src_storage = vkt.CreateMappedStorage<int>(count, VK_BUFFER_USAGE_TRANSFER_SRC_BIT);
|
||||
const auto dst_storage = vkt.CreateMappedStorage<int>(count, VK_BUFFER_USAGE_TRANSFER_DST_BIT);
|
||||
const auto src_storage = vkt.CreateMappedStorage<int>(count,
|
||||
VK_BUFFER_USAGE_TRANSFER_SRC_BIT);
|
||||
const auto dst_storage = vkt.CreateMappedStorage<int>(count,
|
||||
VK_BUFFER_USAGE_TRANSFER_DST_BIT);
|
||||
|
||||
#if HT_AMD
|
||||
constexpr auto second_semaphore_type = VK_SEMAPHORE_TYPE_BINARY;
|
||||
@@ -128,7 +128,8 @@ TEST_CASE("Unit_hipSignalExternalSemaphoresAsync_Vulkan_Positive_Multiple_Semaph
|
||||
const auto hip_timeline_sem_handle_desc =
|
||||
vkt.BuildSemaphoreDescriptor(timeline_semaphore, second_semaphore_type);
|
||||
hipExternalSemaphore_t hip_timeline_ext_semaphore;
|
||||
HIP_CHECK(hipImportExternalSemaphore(&hip_timeline_ext_semaphore, &hip_timeline_sem_handle_desc));
|
||||
HIP_CHECK(hipImportExternalSemaphore(&hip_timeline_ext_semaphore,
|
||||
&hip_timeline_sem_handle_desc));
|
||||
|
||||
uint64_t wait_values[] = {1, 0};
|
||||
VkTimelineSemaphoreSubmitInfo timeline_submit_info = {};
|
||||
@@ -156,7 +157,8 @@ TEST_CASE("Unit_hipSignalExternalSemaphoresAsync_Vulkan_Positive_Multiple_Semaph
|
||||
timeline_signal_params.params.fence.value =
|
||||
second_semaphore_type == VK_SEMAPHORE_TYPE_TIMELINE ? 2 : 0;
|
||||
hipExternalSemaphore_t ext_semaphores[] = {hip_binary_ext_semaphore, hip_timeline_ext_semaphore};
|
||||
hipExternalSemaphoreSignalParams signal_params[] = {binary_signal_params, timeline_signal_params};
|
||||
hipExternalSemaphoreSignalParams signal_params[] = {binary_signal_params,
|
||||
timeline_signal_params};
|
||||
HIP_CHECK(hipSignalExternalSemaphoresAsync(ext_semaphores, signal_params, 2, nullptr));
|
||||
|
||||
VK_CHECK_RESULT(
|
||||
@@ -165,6 +167,7 @@ TEST_CASE("Unit_hipSignalExternalSemaphoresAsync_Vulkan_Positive_Multiple_Semaph
|
||||
HIP_CHECK(hipDestroyExternalSemaphore(hip_binary_ext_semaphore));
|
||||
HIP_CHECK(hipDestroyExternalSemaphore(hip_timeline_ext_semaphore));
|
||||
}
|
||||
#endif
|
||||
|
||||
TEST_CASE("Unit_hipSignalExternalSemaphoresAsync_Vulkan_Negative_Parameters") {
|
||||
VulkanTest vkt(enable_validation);
|
||||
@@ -197,8 +200,8 @@ TEST_CASE("Unit_hipSignalExternalSemaphoresAsync_Vulkan_Negative_Parameters") {
|
||||
hipStream_t stream = nullptr;
|
||||
HIP_CHECK(hipStreamCreate(&stream));
|
||||
HIP_CHECK(hipStreamDestroy(stream));
|
||||
HIP_CHECK_ERROR(hipSignalExternalSemaphoresAsync(&hip_ext_semaphore, &signal_params, 1, stream),
|
||||
hipErrorInvalidValue);
|
||||
HIP_CHECK_ERROR(hipSignalExternalSemaphoresAsync(&hip_ext_semaphore, &signal_params, 1,
|
||||
stream), hipErrorInvalidValue);
|
||||
HIP_CHECK(hipDestroyExternalSemaphore(hip_ext_semaphore));
|
||||
}
|
||||
}
|
||||
@@ -39,11 +39,12 @@ VkFence VulkanTest::CreateFence() {
|
||||
|
||||
VkSemaphore VulkanTest::CreateExternalSemaphore(VkSemaphoreType sem_type, uint64_t initial_value) {
|
||||
VkExportSemaphoreCreateInfoKHR export_sem_create_info = {};
|
||||
VkSemaphoreTypeCreateInfo timeline_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;
|
||||
|
||||
@@ -33,7 +33,6 @@ THE SOFTWARE.
|
||||
#include <aclapi.h>
|
||||
#include <securitybaseapi.h>
|
||||
#include <vulkan/vulkan_win32.h>
|
||||
|
||||
#endif
|
||||
|
||||
#include <vector>
|
||||
@@ -58,7 +57,7 @@ protected:
|
||||
PSECURITY_DESCRIPTOR m_winPSecurityDescriptor;
|
||||
|
||||
public:
|
||||
WindowsSecurityAttributes::WindowsSecurityAttributes()
|
||||
WindowsSecurityAttributes()
|
||||
{
|
||||
m_winPSecurityDescriptor = (PSECURITY_DESCRIPTOR)calloc(1, SECURITY_DESCRIPTOR_MIN_LENGTH + 2 * sizeof(void **));
|
||||
if (!m_winPSecurityDescriptor) {
|
||||
@@ -92,12 +91,12 @@ WindowsSecurityAttributes::WindowsSecurityAttributes()
|
||||
}
|
||||
|
||||
SECURITY_ATTRIBUTES *
|
||||
WindowsSecurityAttributes::operator&()
|
||||
operator&()
|
||||
{
|
||||
return &m_winSecurityAttributes;
|
||||
}
|
||||
|
||||
WindowsSecurityAttributes::~WindowsSecurityAttributes()
|
||||
~WindowsSecurityAttributes()
|
||||
{
|
||||
PSID *ppSID = (PSID *)((PBYTE)m_winPSecurityDescriptor + SECURITY_DESCRIPTOR_MIN_LENGTH);
|
||||
PACL *ppACL = (PACL *)((PBYTE)ppSID + sizeof(PSID *));
|
||||
@@ -292,7 +291,13 @@ VulkanTest::MappedBuffer<T> VulkanTest::CreateMappedStorage(uint32_t count,
|
||||
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);
|
||||
if (allocate_info.memoryTypeIndex == VK_MAX_MEMORY_TYPES) {
|
||||
WARN("Not supported memory type "
|
||||
<< memory_requirements.memoryTypeBits
|
||||
<<
|
||||
" with VK_MEMORY_PROPERTY_HOST_COHERENT_BIT | VK_MEMORY_PROPERTY_HOST_VISIBLE_BIT");
|
||||
return MappedBuffer<T>{nullptr, nullptr, 0, nullptr};
|
||||
}
|
||||
|
||||
VkExportMemoryAllocateInfoKHR vulkan_export_memory_allocate_info = {};
|
||||
if (external) {
|
||||
@@ -333,4 +338,4 @@ VulkanTest::MappedBuffer<T> VulkanTest::CreateMappedStorage(uint32_t count,
|
||||
// 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);
|
||||
hipExternalSemaphore_t ImportBinarySemaphore(VulkanTest& vkt);
|
||||
@@ -1,18 +1,19 @@
|
||||
Windows
|
||||
--------
|
||||
Prepare
|
||||
• Install hip and visual studio
|
||||
• Install vulkan sdk from vulkan.lunarg.com
|
||||
• Download GLFW binaries from glfw.org
|
||||
• Convert sinwave.farg and vert files to spv
|
||||
o c:\VulkanSDK\1.2.182.0\bin\glslangValidator.exe sinewave.vert -V -o vert.spv
|
||||
o c:\VulkanSDK\1.2.182.0\bin\glslangValidator.exe sinewave.frag -V -o frag.spv
|
||||
• c:\VulkanSDK\1.3.243.0\bin\glslangValidator.exe sinewave.vert -V -o vert.spv
|
||||
• c:\VulkanSDK\1.3.243.0\bin\glslangValidator.exe sinewave.frag -V -o frag.spv
|
||||
|
||||
to build without cmake:
|
||||
Build without CMake
|
||||
• set HCC_AMDGPU_TARGET=gfx906:sramecc-:xnack- (for your graphic card, you can get the name from hipinfo )
|
||||
• hipcc -v *.cpp *.hip -Ic:\VulkanSDK\1.2.182.0\include -L c:\VulkanSDK\1.2.182.0\lib -Ic:\glfw-3.3.4.bin.WIN64\include -L c:\glfw-3.3.4.bin.WIN64\lib-vc2019 -Ic:\hip\include\hip -lglfw3dll -lvulkan-1 -ladvapi32 -std=c++14
|
||||
• run a.exe, you should see a 3D sinewave simulation
|
||||
• hipcc -v *.cpp *.hip -o hip_vulkan_image.exe -Ic:\VulkanSDK\1.3.243.0\include -L c:\VulkanSDK\1.3.243.0\lib -Ic:\glfw-3.3.8.bin.WIN64\include -L c:\glfw-3.3.8.bin.WIN64\lib-vc2019 -Ic:\hip\include\hip -lglfw3dll -lvulkan-1 -ladvapi32 -std=c++14
|
||||
• run hip_vulkan_image.exe, you should see a 3D sinewave simulation
|
||||
|
||||
to build with cmake on windows:
|
||||
Build with CMake:
|
||||
• mkdir build; cd build
|
||||
• cmake.exe -GNinja -DCMAKE_CXX_COMPILER_ID=ROCMClang -DCMAKE_C_COMPILER_ID=ROCMClang -DCMAKE_PREFIX_PATH=d:\driver2\drivers\drivers\compute\hip_sdk
|
||||
|
||||
@@ -27,4 +28,3 @@ Build with CMake:
|
||||
• cmake -DCMAKE_PREFIX_PATH=path\to\rocm -DHIP_CXX_COMPILER=path\to\clang
|
||||
• make
|
||||
• run hipVulkan executable
|
||||
|
||||
|
||||
مرجع در شماره جدید
Block a user