diff --git a/test_common/harness/errorHelpers.h b/test_common/harness/errorHelpers.h index c302397eaf..3f1d8fb145 100644 --- a/test_common/harness/errorHelpers.h +++ b/test_common/harness/errorHelpers.h @@ -61,6 +61,21 @@ static int vlog_win32(const char *format, ...); log_error(msg, ##__VA_ARGS__); \ return TEST_FAIL; \ } +#define test_fail_and_cleanup(errRet, cleanup, msg, ...) \ + { \ + log_error(msg, ##__VA_ARGS__); \ + errRet = TEST_FAIL; \ + goto cleanup; \ + } +#define test_error_and_cleanup(errCode, cleanup, msg, ...) \ + { \ + auto errCodeResult = errCode; \ + if (errCodeResult != CL_SUCCESS) \ + { \ + print_error(errCodeResult, msg); \ + goto cleanup; \ + } \ + } #define test_error(errCode, msg) test_error_ret(errCode, msg, errCode) #define test_error_fail(errCode, msg) test_error_ret(errCode, msg, TEST_FAIL) #define test_error_ret(errCode, msg, retValue) \ diff --git a/test_common/harness/fpcontrol.h b/test_common/harness/fpcontrol.h index 222aa2c40c..12aba0a94a 100644 --- a/test_common/harness/fpcontrol.h +++ b/test_common/harness/fpcontrol.h @@ -43,6 +43,8 @@ typedef int64_t FPU_mode_type; #elif defined(__PPC__) #include extern __thread fpu_control_t fpu_control; +#elif defined(__mips__) +#include "mips/m32c1.h" #endif // Set the reference hardware floating point unit to FTZ mode inline void ForceFTZ(FPU_mode_type *mode) @@ -65,6 +67,8 @@ inline void ForceFTZ(FPU_mode_type *mode) __asm__ volatile("mrs %0, fpcr" : "=r"(fpscr)); *mode = fpscr; __asm__ volatile("msr fpcr, %0" ::"r"(fpscr | (1U << 24))); +#elif defined(__mips__) + fpa_bissr(FPA_CSR_FS); #else #error ForceFTZ needs an implentation #endif @@ -91,6 +95,8 @@ inline void DisableFTZ(FPU_mode_type *mode) __asm__ volatile("mrs %0, fpcr" : "=r"(fpscr)); *mode = fpscr; __asm__ volatile("msr fpcr, %0" ::"r"(fpscr & ~(1U << 24))); +#elif defined(__mips__) + fpa_bicsr(FPA_CSR_FS); #else #error DisableFTZ needs an implentation #endif @@ -109,6 +115,8 @@ inline void RestoreFPState(FPU_mode_type *mode) // Add 64 bit support #elif defined(__aarch64__) __asm__ volatile("msr fpcr, %0" ::"r"(*mode)); +#elif defined(__mips__) + // Mips runs by default with DAZ=1 FTZ=1 #else #error RestoreFPState needs an implementation #endif diff --git a/test_common/harness/rounding_mode.cpp b/test_common/harness/rounding_mode.cpp index 1f531478cf..191c04d929 100644 --- a/test_common/harness/rounding_mode.cpp +++ b/test_common/harness/rounding_mode.cpp @@ -197,6 +197,8 @@ RoundingMode get_round(void) #include #elif defined(__PPC__) #include +#elif defined(__mips__) +#include "mips/m32c1.h" #endif void *FlushToZero(void) { @@ -219,6 +221,9 @@ void *FlushToZero(void) flags |= _FPU_MASK_NI; _FPU_SETCW(flags); return NULL; +#elif defined(__mips__) + fpa_bissr(FPA_CSR_FS); + return NULL; #else #error Unknown arch #endif @@ -247,6 +252,8 @@ void UnFlushToZero(void *p) _FPU_GETCW(flags); flags &= ~_FPU_MASK_NI; _FPU_SETCW(flags); +#elif defined(__mips__) + fpa_bicsr(FPA_CSR_FS); #else #error Unknown arch #endif diff --git a/test_common/harness/testHarness.cpp b/test_common/harness/testHarness.cpp index 3d743e717f..3e5d7c9501 100644 --- a/test_common/harness/testHarness.cpp +++ b/test_common/harness/testHarness.cpp @@ -1298,6 +1298,8 @@ void PrintArch(void) vlog("ARCH:\taarch64\n"); #elif defined(_WIN32) vlog("ARCH:\tWindows\n"); +#elif defined(__mips__) + vlog("ARCH:\tmips\n"); #else #error unknown arch #endif diff --git a/test_conformance/allocations/allocation_execute.cpp b/test_conformance/allocations/allocation_execute.cpp index 5a77c3a75d..fb19cccc73 100644 --- a/test_conformance/allocations/allocation_execute.cpp +++ b/test_conformance/allocations/allocation_execute.cpp @@ -140,7 +140,7 @@ int check_image(cl_command_queue queue, cl_mem mem) { } -#define NUM_OF_WORK_ITEMS 8192*2 +#define NUM_OF_WORK_ITEMS (8192 * 32) int execute_kernel(cl_context context, cl_command_queue *queue, cl_device_id device_id, int test, cl_mem mems[], int number_of_mems_used, int verify_checksum) { diff --git a/test_conformance/c11_atomics/common.h b/test_conformance/c11_atomics/common.h index 37c37e874f..76d1fe2742 100644 --- a/test_conformance/c11_atomics/common.h +++ b/test_conformance/c11_atomics/common.h @@ -644,8 +644,8 @@ class CBasicTestMemOrderScope } private: - TExplicitMemoryOrderType _memoryOrder; - TExplicitMemoryScopeType _memoryScope; + TExplicitMemoryOrderType _memoryOrder = MEMORY_ORDER_EMPTY; + TExplicitMemoryScopeType _memoryScope = MEMORY_SCOPE_EMPTY; }; template diff --git a/test_conformance/c11_atomics/host_atomics.h b/test_conformance/c11_atomics/host_atomics.h index b865970f44..efa36ad304 100644 --- a/test_conformance/c11_atomics/host_atomics.h +++ b/test_conformance/c11_atomics/host_atomics.h @@ -53,12 +53,12 @@ enum TExplicitMemoryOrderType #endif #define HOST_ATOMIC_INTPTR_T32 HOST_ATOMIC_INT -#define HOST_ATOMIC_UINTPTR_T32 HOST_ATOMIC_INT +#define HOST_ATOMIC_UINTPTR_T32 HOST_ATOMIC_UINT #define HOST_ATOMIC_SIZE_T32 HOST_ATOMIC_UINT #define HOST_ATOMIC_PTRDIFF_T32 HOST_ATOMIC_INT #define HOST_ATOMIC_INTPTR_T64 HOST_ATOMIC_LONG -#define HOST_ATOMIC_UINTPTR_T64 HOST_ATOMIC_LONG +#define HOST_ATOMIC_UINTPTR_T64 HOST_ATOMIC_ULONG #define HOST_ATOMIC_SIZE_T64 HOST_ATOMIC_ULONG #define HOST_ATOMIC_PTRDIFF_T64 HOST_ATOMIC_LONG @@ -82,7 +82,7 @@ enum TExplicitMemoryOrderType #define HOST_SIZE_T64 cl_ulong #define HOST_PTRDIFF_T64 cl_long -#define HOST_FLAG cl_uint +#define HOST_FLAG cl_int // host atomic functions void host_atomic_thread_fence(TExplicitMemoryOrderType order); diff --git a/test_conformance/common/vulkan_wrapper/opencl_vulkan_wrapper.cpp b/test_conformance/common/vulkan_wrapper/opencl_vulkan_wrapper.cpp index 5d0e99e0f1..b69be1197c 100644 --- a/test_conformance/common/vulkan_wrapper/opencl_vulkan_wrapper.cpp +++ b/test_conformance/common/vulkan_wrapper/opencl_vulkan_wrapper.cpp @@ -33,6 +33,7 @@ pfnclEnqueueAcquireExternalMemObjectsKHR pfnclEnqueueReleaseExternalMemObjectsKHR clEnqueueReleaseExternalMemObjectsKHRptr; pfnclReleaseSemaphoreKHR clReleaseSemaphoreKHRptr; +pfnclGetSemaphoreHandleForTypeKHR clGetSemaphoreHandleForTypeKHRptr; void init_cl_vk_ext(cl_platform_id opencl_platform) { @@ -69,6 +70,15 @@ void init_cl_vk_ext(cl_platform_id opencl_platform) throw std::runtime_error("Failed to get the function pointer of " "clCreateSemaphoreWithPropertiesKHRptr!"); } + + clGetSemaphoreHandleForTypeKHRptr = (pfnclGetSemaphoreHandleForTypeKHR) + clGetExtensionFunctionAddressForPlatform( + opencl_platform, "clGetSemaphoreHandleForTypeKHR"); + if (NULL == clGetSemaphoreHandleForTypeKHRptr) + { + throw std::runtime_error("Failed to get the function pointer of " + "clGetSemaphoreHandleForTypeKHRptr!"); + } } cl_int setMaxImageDimensions(cl_device_id deviceID, size_t &max_width, @@ -522,8 +532,8 @@ clExternalMemory::clExternalMemory(const clExternalMemory &externalMemory) clExternalMemory::clExternalMemory( const VulkanDeviceMemory *deviceMemory, - VulkanExternalMemoryHandleType externalMemoryHandleType, uint64_t offset, - uint64_t size, cl_context context, cl_device_id deviceId) + VulkanExternalMemoryHandleType externalMemoryHandleType, uint64_t size, + cl_context context, cl_device_id deviceId) { int err = 0; m_externalMemory = NULL; @@ -548,9 +558,9 @@ clExternalMemory::clExternalMemory( { case VULKAN_EXTERNAL_MEMORY_HANDLE_TYPE_OPAQUE_FD: #ifdef _WIN32 + log_info("Opaque file descriptors are not supported on Windows\n"); ASSERT(0); #endif - log_info("Opaque file descriptors are not supported on Windows\n"); fd = (int)deviceMemory->getHandle(externalMemoryHandleType); err = check_external_memory_handle_type( devList[0], CL_EXTERNAL_MEMORY_HANDLE_OPAQUE_FD_KHR); @@ -595,10 +605,11 @@ clExternalMemory::clExternalMemory( throw std::runtime_error("Unsupported external memory type\n "); } - extMemProperties.push_back((cl_mem_properties)CL_DEVICE_HANDLE_LIST_KHR); + extMemProperties.push_back( + (cl_mem_properties)CL_MEM_DEVICE_HANDLE_LIST_KHR); extMemProperties.push_back((cl_mem_properties)devList[0]); extMemProperties.push_back( - (cl_mem_properties)CL_DEVICE_HANDLE_LIST_END_KHR); + (cl_mem_properties)CL_MEM_DEVICE_HANDLE_LIST_END_KHR); extMemProperties.push_back(0); m_externalMemory = clCreateBufferWithProperties( @@ -691,10 +702,11 @@ clExternalMemoryImage::clExternalMemoryImage( throw std::runtime_error("getCLImageInfoFromVkImageInfo failed!!!"); } - extMemProperties1.push_back((cl_mem_properties)CL_DEVICE_HANDLE_LIST_KHR); + extMemProperties1.push_back( + (cl_mem_properties)CL_MEM_DEVICE_HANDLE_LIST_KHR); extMemProperties1.push_back((cl_mem_properties)devList[0]); extMemProperties1.push_back( - (cl_mem_properties)CL_DEVICE_HANDLE_LIST_END_KHR); + (cl_mem_properties)CL_MEM_DEVICE_HANDLE_LIST_END_KHR); extMemProperties1.push_back(0); m_externalMemory = clCreateImageWithProperties( context, extMemProperties1.data(), CL_MEM_READ_WRITE, &img_format, @@ -726,56 +738,19 @@ clExternalMemoryImage::clExternalMemoryImage() {} // clExternalSemaphore implementation // ////////////////////////////////////////// -clExternalSemaphore::clExternalSemaphore( - const clExternalSemaphore &externalSemaphore) - : m_externalSemaphore(externalSemaphore.m_externalSemaphore) -{} - clExternalSemaphore::clExternalSemaphore( const VulkanSemaphore &semaphore, cl_context context, VulkanExternalSemaphoreHandleType externalSemaphoreHandleType, cl_device_id deviceId) + : m_deviceSemaphore(semaphore) { cl_int err = 0; cl_device_id devList[] = { deviceId, NULL }; - - switch (externalSemaphoreHandleType) - { - case VULKAN_EXTERNAL_SEMAPHORE_HANDLE_TYPE_OPAQUE_FD: - if (!is_extension_available(devList[0], - "cl_khr_external_semaphore_opaque_fd")) - { - throw std::runtime_error("Device does not support " - "cl_khr_external_semaphore_opaque_fd " - "extension \n"); - } - break; - case VULKAN_EXTERNAL_SEMAPHORE_HANDLE_TYPE_OPAQUE_WIN32_NT: - case VULKAN_EXTERNAL_SEMAPHORE_HANDLE_TYPE_OPAQUE_WIN32_KMT: - case VULKAN_EXTERNAL_SEMAPHORE_HANDLE_TYPE_OPAQUE_WIN32_NT_KMT: - if (!is_extension_available(devList[0], - "cl_khr_external_semaphore_win32")) - { - throw std::runtime_error( - "Device does not support " - "cl_khr_external_semaphore_win32 extension\n"); - } - break; - case VULKAN_EXTERNAL_SEMAPHORE_HANDLE_TYPE_SYNC_FD: - if (!is_extension_available(devList[0], - "cl_khr_external_semaphore_sync_fd")) - { - throw std::runtime_error( - "Device does not support cl_khr_external_semaphore_sync_fd " - "extension \n"); - } - break; - default: - throw std::runtime_error( - "Unsupported external semaphore handle type\n"); - break; - } + m_externalHandleType = externalSemaphoreHandleType; + m_externalSemaphore = nullptr; + m_device = deviceId; + m_context = context; std::vector sema_props{ (cl_semaphore_properties_khr)CL_SEMAPHORE_TYPE_KHR, @@ -846,10 +821,10 @@ clExternalSemaphore::clExternalSemaphore( } sema_props.push_back( - (cl_semaphore_properties_khr)CL_DEVICE_HANDLE_LIST_KHR); + (cl_semaphore_properties_khr)CL_SEMAPHORE_DEVICE_HANDLE_LIST_KHR); sema_props.push_back((cl_semaphore_properties_khr)devList[0]); sema_props.push_back( - (cl_semaphore_properties_khr)CL_DEVICE_HANDLE_LIST_END_KHR); + (cl_semaphore_properties_khr)CL_SEMAPHORE_DEVICE_HANDLE_LIST_END_KHR); sema_props.push_back(0); m_externalSemaphore = clCreateSemaphoreWithPropertiesKHRptr(context, sema_props.data(), &err); @@ -871,16 +846,97 @@ clExternalSemaphore::~clExternalSemaphore() noexcept(false) } } -void clExternalSemaphore::signal(cl_command_queue cmd_queue) +int clExternalSemaphore::signal(cl_command_queue cmd_queue) { - clEnqueueSignalSemaphoresKHRptr(cmd_queue, 1, &m_externalSemaphore, NULL, 0, - NULL, NULL); + int err = clEnqueueSignalSemaphoresKHRptr( + cmd_queue, 1, &m_externalSemaphore, NULL, 0, NULL, nullptr); + if (err != CL_SUCCESS) + { + return err; + } + + if (m_externalHandleType == VULKAN_EXTERNAL_SEMAPHORE_HANDLE_TYPE_SYNC_FD) + { + err = clGetSemaphoreHandleForTypeKHRptr(m_externalSemaphore, m_device, + CL_SEMAPHORE_HANDLE_SYNC_FD_KHR, + sizeof(int), &fd, nullptr); + if (err != CL_SUCCESS) + { + log_error("Failed to export fd from semaphore\n"); + return err; + } + + VkImportSemaphoreFdInfoKHR import = {}; + import.sType = VK_STRUCTURE_TYPE_IMPORT_SEMAPHORE_FD_INFO_KHR; + import.semaphore = m_deviceSemaphore; + import.fd = fd; + import.pNext = nullptr; + import.handleType = VK_EXTERNAL_SEMAPHORE_HANDLE_TYPE_SYNC_FD_BIT_KHR; + import.flags = 0; + + VkResult res = + vkImportSemaphoreFdKHR(m_deviceSemaphore.getDevice(), &import); + ASSERT(res == VK_SUCCESS); + if (res != VK_SUCCESS) + { + err = CL_INVALID_OPERATION; + } + } + + return err; } -void clExternalSemaphore::wait(cl_command_queue cmd_queue) +int clExternalSemaphore::wait(cl_command_queue cmd_queue) { - clEnqueueWaitSemaphoresKHRptr(cmd_queue, 1, &m_externalSemaphore, NULL, 0, - NULL, NULL); + int err = CL_SUCCESS; + if (m_externalHandleType == VULKAN_EXTERNAL_SEMAPHORE_HANDLE_TYPE_SYNC_FD) + { + cl_int err = 0; + cl_device_id devList[] = { m_device, NULL }; + std::vector sema_props{ + (cl_semaphore_properties_khr)CL_SEMAPHORE_TYPE_KHR, + (cl_semaphore_properties_khr)CL_SEMAPHORE_TYPE_BINARY_KHR, + }; + fd = (int)m_deviceSemaphore.getHandle(m_externalHandleType); + + err = check_external_semaphore_handle_type( + devList[0], CL_SEMAPHORE_HANDLE_SYNC_FD_KHR); + if (CL_SUCCESS != err) + { + log_error("CL_SEMAPHORE_HANDLE_SYNC_FD_KHR not supported\n"); + return err; + } + + sema_props.push_back( + (cl_semaphore_properties_khr)CL_SEMAPHORE_HANDLE_SYNC_FD_KHR); + sema_props.push_back((cl_semaphore_properties_khr)fd); + + sema_props.push_back(0); + + if (m_externalSemaphore) + { + err = clReleaseSemaphoreKHRptr(m_externalSemaphore); + if (err != CL_SUCCESS) + { + log_error("Failed to release CL external semaphore\n"); + return err; + } + m_externalSemaphore = nullptr; + } + + m_externalSemaphore = clCreateSemaphoreWithPropertiesKHRptr( + m_context, sema_props.data(), &err); + if (CL_SUCCESS != err) + { + log_error("clCreateSemaphoreWithPropertiesKHRptr failed with %d\n", + err); + return err; + } + } + + err = clEnqueueWaitSemaphoresKHRptr(cmd_queue, 1, &m_externalSemaphore, + NULL, 0, NULL, NULL); + return err; } cl_semaphore_khr &clExternalSemaphore::getCLSemaphore() diff --git a/test_conformance/common/vulkan_wrapper/opencl_vulkan_wrapper.hpp b/test_conformance/common/vulkan_wrapper/opencl_vulkan_wrapper.hpp index 4a1d453e59..12d467d8be 100644 --- a/test_conformance/common/vulkan_wrapper/opencl_vulkan_wrapper.hpp +++ b/test_conformance/common/vulkan_wrapper/opencl_vulkan_wrapper.hpp @@ -50,6 +50,10 @@ typedef cl_int (*pfnclEnqueueReleaseExternalMemObjectsKHR)( const cl_mem *mem_objects, cl_uint num_events_in_wait_list, const cl_event *event_wait_list, cl_event *event); typedef cl_int (*pfnclReleaseSemaphoreKHR)(cl_semaphore_khr sema_object); +typedef cl_int (*pfnclGetSemaphoreHandleForTypeKHR)( + cl_semaphore_khr sema_object, cl_device_id device, + cl_external_semaphore_handle_type_khr handleType, size_t handle_size, + void *handle, size_t *handleSize); extern pfnclCreateSemaphoreWithPropertiesKHR clCreateSemaphoreWithPropertiesKHRptr; @@ -83,8 +87,7 @@ class clExternalMemory { clExternalMemory(); clExternalMemory(const VulkanDeviceMemory *deviceMemory, VulkanExternalMemoryHandleType externalMemoryHandleType, - uint64_t offset, uint64_t size, cl_context context, - cl_device_id deviceId); + uint64_t size, cl_context context, cl_device_id deviceId); virtual ~clExternalMemory(); cl_mem getExternalMemoryBuffer(); @@ -111,9 +114,12 @@ class clExternalMemoryImage { class clExternalSemaphore { protected: cl_semaphore_khr m_externalSemaphore; + VulkanExternalSemaphoreHandleType m_externalHandleType; + cl_device_id m_device; + cl_context m_context; + const VulkanSemaphore &m_deviceSemaphore; int fd; void *handle; - clExternalSemaphore(const clExternalSemaphore &externalSemaphore); public: clExternalSemaphore( @@ -121,8 +127,8 @@ class clExternalSemaphore { VulkanExternalSemaphoreHandleType externalSemaphoreHandleType, cl_device_id deviceId); virtual ~clExternalSemaphore() noexcept(false); - void signal(cl_command_queue command_queue); - void wait(cl_command_queue command_queue); + int signal(cl_command_queue command_queue); + int wait(cl_command_queue command_queue); cl_semaphore_khr &getCLSemaphore(); // operator openclExternalSemaphore_t() const; }; diff --git a/test_conformance/common/vulkan_wrapper/vulkan_api_list.hpp b/test_conformance/common/vulkan_wrapper/vulkan_api_list.hpp index e9c06f9854..70c0944ed4 100644 --- a/test_conformance/common/vulkan_wrapper/vulkan_api_list.hpp +++ b/test_conformance/common/vulkan_wrapper/vulkan_api_list.hpp @@ -98,8 +98,9 @@ VK_FUNC_DECL(vkGetPhysicalDeviceSurfaceFormatsKHR) \ VK_FUNC_DECL(vkGetPhysicalDeviceSurfacePresentModesKHR) \ VK_FUNC_DECL(vkEnumerateDeviceExtensionProperties) \ - VK_FUNC_DECL(vkGetPhysicalDeviceSurfaceSupportKHR) - + VK_FUNC_DECL(vkGetPhysicalDeviceSurfaceSupportKHR) \ + VK_FUNC_DECL(vkImportSemaphoreFdKHR) \ + VK_FUNC_DECL(vkGetPhysicalDeviceExternalSemaphorePropertiesKHR) #define VK_WINDOWS_FUNC_LIST \ VK_FUNC_DECL(vkGetMemoryWin32HandleKHR) \ VK_FUNC_DECL(vkGetSemaphoreWin32HandleKHR) @@ -192,7 +193,9 @@ _vkEnumerateDeviceExtensionProperties #define vkGetPhysicalDeviceSurfaceSupportKHR \ _vkGetPhysicalDeviceSurfaceSupportKHR - +#define vkImportSemaphoreFdKHR _vkImportSemaphoreFdKHR +#define vkGetPhysicalDeviceExternalSemaphorePropertiesKHR \ + _vkGetPhysicalDeviceExternalSemaphorePropertiesKHR #define vkGetMemoryWin32HandleKHR _vkGetMemoryWin32HandleKHR #define vkGetSemaphoreWin32HandleKHR _vkGetSemaphoreWin32HandleKHR diff --git a/test_conformance/common/vulkan_wrapper/vulkan_list_map.hpp b/test_conformance/common/vulkan_wrapper/vulkan_list_map.hpp index ef00b70ac3..7dd099c090 100644 --- a/test_conformance/common/vulkan_wrapper/vulkan_list_map.hpp +++ b/test_conformance/common/vulkan_wrapper/vulkan_list_map.hpp @@ -336,6 +336,8 @@ const VulkanWrapper & return (m_wrapperList.size() > 0) ? m_wrapperList[idx].get() : m_constWrapperList[idx].get(); } + + throw std::runtime_error("Out of bounds operator access"); } template diff --git a/test_conformance/common/vulkan_wrapper/vulkan_utility.cpp b/test_conformance/common/vulkan_wrapper/vulkan_utility.cpp index 2124a275c4..f694954ef9 100644 --- a/test_conformance/common/vulkan_wrapper/vulkan_utility.cpp +++ b/test_conformance/common/vulkan_wrapper/vulkan_utility.cpp @@ -24,6 +24,7 @@ #include #include #include +#include "deviceInfo.h" #if defined(_WIN32) || defined(_WIN64) #include #endif @@ -174,7 +175,7 @@ getVulkanMemoryType(const VulkanDevice &device, } } - // CHECK_LT(mtIdx, memoryTypeList.size()); + ASSERT(mtIdx < memoryTypeList.size()); return memoryTypeList[mtIdx]; } @@ -236,30 +237,112 @@ getSupportedVulkanExternalMemoryHandleTypeList() } const std::vector -getSupportedVulkanExternalSemaphoreHandleTypeList() +getSupportedVulkanExternalSemaphoreHandleTypeList(const VulkanDevice &vkDevice) { + typedef struct + { + const char *extension_name; + VkExternalSemaphoreHandleTypeFlagBits vk_type; + VulkanExternalSemaphoreHandleType enum_type; + } VkSemaphoreHandleMap; + + // Add all known handle types, use Vulkan queries to determine what is + // supported. + std::vector all_known_handle_types; + all_known_handle_types.push_back( + { "VK_KHR_external_semaphore_fd", + VK_EXTERNAL_SEMAPHORE_HANDLE_TYPE_OPAQUE_FD_BIT_KHR, + VULKAN_EXTERNAL_SEMAPHORE_HANDLE_TYPE_OPAQUE_FD }); + all_known_handle_types.push_back( + { "VK_KHR_external_semaphore_fd", + VK_EXTERNAL_SEMAPHORE_HANDLE_TYPE_SYNC_FD_BIT_KHR, + VULKAN_EXTERNAL_SEMAPHORE_HANDLE_TYPE_SYNC_FD }); + all_known_handle_types.push_back( + { "VK_KHR_external_semaphore_win32", + VK_EXTERNAL_SEMAPHORE_HANDLE_TYPE_OPAQUE_WIN32_KMT_BIT_KHR, + VULKAN_EXTERNAL_SEMAPHORE_HANDLE_TYPE_OPAQUE_WIN32_KMT }); + all_known_handle_types.push_back( + { "VK_KHR_external_semaphore_win32", + VK_EXTERNAL_SEMAPHORE_HANDLE_TYPE_OPAQUE_WIN32_BIT_KHR, + VULKAN_EXTERNAL_SEMAPHORE_HANDLE_TYPE_OPAQUE_WIN32_NT }); + std::vector externalSemaphoreHandleTypeList; -#if _WIN32 - if (IsWindows8OrGreater()) + for (auto handle_type : all_known_handle_types) { - externalSemaphoreHandleTypeList.push_back( - VULKAN_EXTERNAL_SEMAPHORE_HANDLE_TYPE_OPAQUE_WIN32_NT); + if (!vkDevice.getPhysicalDevice().hasExtension( + handle_type.extension_name)) + { + continue; + } + + VkPhysicalDeviceExternalSemaphoreInfo handle_query = { + VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_EXTERNAL_SEMAPHORE_INFO, nullptr, + handle_type.vk_type + }; + VkExternalSemaphoreProperties query_result = {}; + vkGetPhysicalDeviceExternalSemaphorePropertiesKHR( + vkDevice.getPhysicalDevice(), &handle_query, &query_result); + if (query_result.externalSemaphoreFeatures + & (VK_EXTERNAL_SEMAPHORE_FEATURE_EXPORTABLE_BIT_KHR + | VK_EXTERNAL_SEMAPHORE_FEATURE_IMPORTABLE_BIT_KHR)) + { + externalSemaphoreHandleTypeList.push_back(handle_type.enum_type); + } } - externalSemaphoreHandleTypeList.push_back( - VULKAN_EXTERNAL_SEMAPHORE_HANDLE_TYPE_OPAQUE_WIN32_KMT); -#elif defined(__ANDROID__) - externalSemaphoreHandleTypeList.push_back( - VULKAN_EXTERNAL_SEMAPHORE_HANDLE_TYPE_SYNC_FD); -#else - externalSemaphoreHandleTypeList.push_back( - VULKAN_EXTERNAL_SEMAPHORE_HANDLE_TYPE_OPAQUE_FD); -#endif return externalSemaphoreHandleTypeList; } +std::vector +getSupportedInteropExternalSemaphoreHandleTypes(cl_device_id device, + VulkanDevice &vkDevice) +{ + const std::vector + supportedVkSemaphoreTypes = + getSupportedVulkanExternalSemaphoreHandleTypeList(vkDevice); + std::vector supportedSemaphoreTypes; + + if (is_extension_available(device, "cl_khr_external_semaphore_opaque_fd") + && std::count(supportedVkSemaphoreTypes.begin(), + supportedVkSemaphoreTypes.end(), + VULKAN_EXTERNAL_SEMAPHORE_HANDLE_TYPE_OPAQUE_FD)) + { + supportedSemaphoreTypes.push_back( + VULKAN_EXTERNAL_SEMAPHORE_HANDLE_TYPE_OPAQUE_FD); + } + + if (is_extension_available(device, "cl_khr_external_semaphore_sync_fd") + && std::count(supportedVkSemaphoreTypes.begin(), + supportedVkSemaphoreTypes.end(), + VULKAN_EXTERNAL_SEMAPHORE_HANDLE_TYPE_SYNC_FD)) + { + supportedSemaphoreTypes.push_back( + VULKAN_EXTERNAL_SEMAPHORE_HANDLE_TYPE_SYNC_FD); + } + + if (is_extension_available(device, "cl_khr_external_semaphore_win32") + && std::count(supportedVkSemaphoreTypes.begin(), + supportedVkSemaphoreTypes.end(), + VULKAN_EXTERNAL_SEMAPHORE_HANDLE_TYPE_OPAQUE_WIN32_NT)) + { + supportedSemaphoreTypes.push_back( + VULKAN_EXTERNAL_SEMAPHORE_HANDLE_TYPE_OPAQUE_WIN32_NT); + } + + if (is_extension_available(device, "cl_khr_external_semaphore_win32") + && std::count(supportedVkSemaphoreTypes.begin(), + supportedVkSemaphoreTypes.end(), + VULKAN_EXTERNAL_SEMAPHORE_HANDLE_TYPE_OPAQUE_WIN32_KMT)) + { + supportedSemaphoreTypes.push_back( + VULKAN_EXTERNAL_SEMAPHORE_HANDLE_TYPE_OPAQUE_WIN32_KMT); + } + + return supportedSemaphoreTypes; +} + const std::vector getSupportedVulkanFormatList() { std::vector formatList; @@ -498,7 +581,6 @@ cl_external_semaphore_handle_type_khr getCLSemaphoreTypeFromVulkanType( clExternalSemaphoreHandleTypeKhr = CL_SEMAPHORE_HANDLE_OPAQUE_WIN32_KHR; break; - case VULKAN_EXTERNAL_SEMAPHORE_HANDLE_TYPE_OPAQUE_WIN32_NT_KMT: case VULKAN_EXTERNAL_SEMAPHORE_HANDLE_TYPE_OPAQUE_WIN32_KMT: clExternalSemaphoreHandleTypeKhr = CL_SEMAPHORE_HANDLE_OPAQUE_WIN32_KMT_KHR; @@ -631,8 +713,8 @@ operator<<(std::ostream &os, return os << "Opaque NT handle"; case VULKAN_EXTERNAL_SEMAPHORE_HANDLE_TYPE_OPAQUE_WIN32_KMT: return os << "Opaque D3DKMT handle"; - case VULKAN_EXTERNAL_SEMAPHORE_HANDLE_TYPE_OPAQUE_WIN32_NT_KMT: - return os << "Opaque NT and D3DKMT handle"; + case VULKAN_EXTERNAL_SEMAPHORE_HANDLE_TYPE_SYNC_FD: + return os << "Sync fd semaphore handle"; } return os; diff --git a/test_conformance/common/vulkan_wrapper/vulkan_utility.hpp b/test_conformance/common/vulkan_wrapper/vulkan_utility.hpp index 51284125b4..b3cc7e44cc 100644 --- a/test_conformance/common/vulkan_wrapper/vulkan_utility.hpp +++ b/test_conformance/common/vulkan_wrapper/vulkan_utility.hpp @@ -46,7 +46,10 @@ getDefaultVulkanQueueFamilyToQueueCountMap(); const std::vector getSupportedVulkanExternalMemoryHandleTypeList(); const std::vector -getSupportedVulkanExternalSemaphoreHandleTypeList(); +getSupportedVulkanExternalSemaphoreHandleTypeList(const VulkanDevice& vkDevice); +std::vector +getSupportedInteropExternalSemaphoreHandleTypes(cl_device_id device, + VulkanDevice& vkDevice); const std::vector getSupportedVulkanFormatList(); uint32_t getVulkanFormatElementSize(VulkanFormat format); diff --git a/test_conformance/common/vulkan_wrapper/vulkan_wrapper.cpp b/test_conformance/common/vulkan_wrapper/vulkan_wrapper.cpp index 73c5e9a13d..b187181f3f 100644 --- a/test_conformance/common/vulkan_wrapper/vulkan_wrapper.cpp +++ b/test_conformance/common/vulkan_wrapper/vulkan_wrapper.cpp @@ -335,6 +335,16 @@ VulkanPhysicalDevice::VulkanPhysicalDevice(VkPhysicalDevice vkPhysicalDevice) memoryHeap); m_memoryTypeList.add(*memoryType); } + + uint32_t num_extensions = 0; + vkEnumerateDeviceExtensionProperties(m_vkPhysicalDevice, nullptr, + &num_extensions, nullptr); + if (num_extensions) + { + m_extensions.resize(num_extensions); + vkEnumerateDeviceExtensionProperties( + m_vkPhysicalDevice, nullptr, &num_extensions, m_extensions.data()); + } } VulkanPhysicalDevice::~VulkanPhysicalDevice() @@ -388,6 +398,18 @@ VulkanPhysicalDevice::operator VkPhysicalDevice() const return m_vkPhysicalDevice; } +bool VulkanPhysicalDevice::hasExtension(const char *extension_name) const +{ + for (const auto &m_extension : m_extensions) + { + if (!strcmp(m_extension.extensionName, extension_name)) + { + return true; + } + } + return false; +} + bool operator<(const VulkanQueueFamily &queueFamilyA, const VulkanQueueFamily &queueFamilyB) { @@ -2256,6 +2278,8 @@ VulkanSemaphore::VulkanSemaphore( vkCreateSemaphore(m_device, &vkSemaphoreCreateInfo, NULL, &m_vkSemaphore); } +const VulkanDevice &VulkanSemaphore::getDevice() const { return m_device; } + VulkanSemaphore::~VulkanSemaphore() { vkDestroySemaphore(m_device, m_vkSemaphore, NULL); @@ -2301,6 +2325,23 @@ int VulkanSemaphore::getHandle( return fd; } + else if (externalSemaphoreHandleType + == VULKAN_EXTERNAL_SEMAPHORE_HANDLE_TYPE_SYNC_FD) + { + int fd; + + VkSemaphoreGetFdInfoKHR vkSemaphoreGetFdInfoKHR = {}; + vkSemaphoreGetFdInfoKHR.sType = + VK_STRUCTURE_TYPE_SEMAPHORE_GET_FD_INFO_KHR; + vkSemaphoreGetFdInfoKHR.pNext = NULL; + vkSemaphoreGetFdInfoKHR.semaphore = m_vkSemaphore; + vkSemaphoreGetFdInfoKHR.handleType = + VK_EXTERNAL_SEMAPHORE_HANDLE_TYPE_SYNC_FD_BIT_KHR; + + vkGetSemaphoreFdKHR(m_device, &vkSemaphoreGetFdInfoKHR, &fd); + + return fd; + } return HANDLE_ERROR; } #endif diff --git a/test_conformance/common/vulkan_wrapper/vulkan_wrapper.hpp b/test_conformance/common/vulkan_wrapper/vulkan_wrapper.hpp index 7fcc70f374..a520dceea7 100644 --- a/test_conformance/common/vulkan_wrapper/vulkan_wrapper.hpp +++ b/test_conformance/common/vulkan_wrapper/vulkan_wrapper.hpp @@ -53,12 +53,15 @@ class VulkanPhysicalDevice { VulkanQueueFamilyList m_queueFamilyList; VulkanMemoryHeapList m_memoryHeapList; VulkanMemoryTypeList m_memoryTypeList; + std::vector m_extensions; + VulkanPhysicalDevice(const VulkanPhysicalDevice &physicalDevice); VulkanPhysicalDevice(VkPhysicalDevice vkPhysicalDevice); virtual ~VulkanPhysicalDevice(); public: + bool hasExtension(const char *extension_name) const; const VulkanQueueFamilyList &getQueueFamilyList() const; const VulkanMemoryHeapList &getMemoryHeapList() const; const VulkanMemoryTypeList &getMemoryTypeList() const; @@ -537,6 +540,7 @@ class VulkanDeviceMemory { uint64_t m_size; bool m_isDedicated; + VulkanDeviceMemory(const VulkanDeviceMemory &deviceMemory); public: @@ -588,6 +592,7 @@ class VulkanSemaphore { VulkanExternalSemaphoreHandleType externalSemaphoreHandleType = VULKAN_EXTERNAL_SEMAPHORE_HANDLE_TYPE_NONE, const std::wstring name = L""); + const VulkanDevice &getDevice() const; virtual ~VulkanSemaphore(); #ifdef _WIN32 HANDLE getHandle( diff --git a/test_conformance/common/vulkan_wrapper/vulkan_wrapper_types.hpp b/test_conformance/common/vulkan_wrapper/vulkan_wrapper_types.hpp index fcd193732f..86d7381322 100644 --- a/test_conformance/common/vulkan_wrapper/vulkan_wrapper_types.hpp +++ b/test_conformance/common/vulkan_wrapper/vulkan_wrapper_types.hpp @@ -167,9 +167,6 @@ enum VulkanExternalSemaphoreHandleType VK_EXTERNAL_SEMAPHORE_HANDLE_TYPE_OPAQUE_WIN32_BIT_KHR, VULKAN_EXTERNAL_SEMAPHORE_HANDLE_TYPE_OPAQUE_WIN32_KMT = VK_EXTERNAL_SEMAPHORE_HANDLE_TYPE_OPAQUE_WIN32_KMT_BIT_KHR, - VULKAN_EXTERNAL_SEMAPHORE_HANDLE_TYPE_OPAQUE_WIN32_NT_KMT = - VK_EXTERNAL_SEMAPHORE_HANDLE_TYPE_OPAQUE_WIN32_BIT_KHR - | VK_EXTERNAL_SEMAPHORE_HANDLE_TYPE_OPAQUE_WIN32_KMT_BIT_KHR, VULKAN_EXTERNAL_SEMAPHORE_HANDLE_TYPE_SYNC_FD = VK_EXTERNAL_SEMAPHORE_HANDLE_TYPE_SYNC_FD_BIT_KHR }; diff --git a/test_conformance/extensions/cl_khr_external_semaphore/test_external_semaphore.cpp b/test_conformance/extensions/cl_khr_external_semaphore/test_external_semaphore.cpp index 89ab17b388..8c0c64f42c 100644 --- a/test_conformance/extensions/cl_khr_external_semaphore/test_external_semaphore.cpp +++ b/test_conformance/extensions/cl_khr_external_semaphore/test_external_semaphore.cpp @@ -54,6 +54,16 @@ static const char* source = "__kernel void empty() {}"; +static void log_info_semaphore_type( + VulkanExternalSemaphoreHandleType vkExternalSemaphoreHandleType) +{ + std::stringstream semaphore_type_description; + semaphore_type_description << "Testing semaphore type \"" + << vkExternalSemaphoreHandleType << "\"" + << std::endl; + log_info("%s", semaphore_type_description.str().c_str()); +} + static int init_vuikan_device() { cl_platform_id platform = nullptr; @@ -104,48 +114,57 @@ int test_external_semaphores_queries(cl_device_id deviceID, cl_context context, GET_PFN(deviceID, clReleaseSemaphoreKHR); GET_PFN(deviceID, clRetainSemaphoreKHR); - const std::vector - vkExternalMemoryHandleTypeList = - getSupportedVulkanExternalMemoryHandleTypeList(); - VulkanExternalSemaphoreHandleType vkExternalSemaphoreHandleType = - getSupportedVulkanExternalSemaphoreHandleTypeList()[0]; - VulkanSemaphore vkVk2CLSemaphore(vkDevice, vkExternalSemaphoreHandleType); + std::vector + vkExternalSemaphoreHandleTypeList = + getSupportedInteropExternalSemaphoreHandleTypes(deviceID, vkDevice); - clExternalSemaphore sema_ext(vkVk2CLSemaphore, context, - vkExternalSemaphoreHandleType, deviceID); + if (vkExternalSemaphoreHandleTypeList.empty()) + { + test_fail("No external semaphore handle types found\n"); + } - // Needed by the macro - cl_semaphore_khr sema = sema_ext.getCLSemaphore(); + for (VulkanExternalSemaphoreHandleType vkExternalSemaphoreHandleType : + vkExternalSemaphoreHandleTypeList) + { + log_info_semaphore_type(vkExternalSemaphoreHandleType); + VulkanSemaphore vkVk2CLSemaphore(vkDevice, + vkExternalSemaphoreHandleType); - SEMAPHORE_PARAM_TEST(CL_SEMAPHORE_TYPE_KHR, cl_semaphore_type_khr, - CL_SEMAPHORE_TYPE_BINARY_KHR); + clExternalSemaphore sema_ext(vkVk2CLSemaphore, context, + vkExternalSemaphoreHandleType, deviceID); - SEMAPHORE_PARAM_TEST(CL_DEVICE_HANDLE_LIST_KHR, cl_device_id, deviceID); + // Needed by the macro + cl_semaphore_khr sema = sema_ext.getCLSemaphore(); - SEMAPHORE_PARAM_TEST( - CL_SEMAPHORE_EXPORT_HANDLE_TYPES_KHR, cl_uint, - getCLSemaphoreTypeFromVulkanType(vkExternalSemaphoreHandleType)); + SEMAPHORE_PARAM_TEST(CL_SEMAPHORE_TYPE_KHR, cl_semaphore_type_khr, + CL_SEMAPHORE_TYPE_BINARY_KHR); - // Confirm that querying CL_SEMAPHORE_CONTEXT_KHR returns the right context - SEMAPHORE_PARAM_TEST(CL_SEMAPHORE_CONTEXT_KHR, cl_context, context); + SEMAPHORE_PARAM_TEST(CL_SEMAPHORE_DEVICE_HANDLE_LIST_KHR, cl_device_id, + deviceID); - // Confirm that querying CL_SEMAPHORE_REFERENCE_COUNT_KHR returns the right - // value - SEMAPHORE_PARAM_TEST(CL_SEMAPHORE_REFERENCE_COUNT_KHR, cl_uint, 1); + // Confirm that querying CL_SEMAPHORE_CONTEXT_KHR returns the right + // context + SEMAPHORE_PARAM_TEST(CL_SEMAPHORE_CONTEXT_KHR, cl_context, context); - cl_int err = CL_SUCCESS; + // Confirm that querying CL_SEMAPHORE_REFERENCE_COUNT_KHR returns the + // right value + SEMAPHORE_PARAM_TEST(CL_SEMAPHORE_REFERENCE_COUNT_KHR, cl_uint, 1); - err = clRetainSemaphoreKHR(sema); - test_error(err, "Could not retain semaphore"); - SEMAPHORE_PARAM_TEST(CL_SEMAPHORE_REFERENCE_COUNT_KHR, cl_uint, 2); + cl_int err = CL_SUCCESS; - err = clReleaseSemaphoreKHR(sema); - test_error(err, "Could not release semaphore"); - SEMAPHORE_PARAM_TEST(CL_SEMAPHORE_REFERENCE_COUNT_KHR, cl_uint, 1); + err = clRetainSemaphoreKHR(sema); + test_error(err, "Could not retain semaphore"); + SEMAPHORE_PARAM_TEST(CL_SEMAPHORE_REFERENCE_COUNT_KHR, cl_uint, 2); - // Confirm that querying CL_SEMAPHORE_PAYLOAD_KHR returns the unsignaled - // state - SEMAPHORE_PARAM_TEST(CL_SEMAPHORE_PAYLOAD_KHR, cl_semaphore_payload_khr, 0); + err = clReleaseSemaphoreKHR(sema); + test_error(err, "Could not release semaphore"); + SEMAPHORE_PARAM_TEST(CL_SEMAPHORE_REFERENCE_COUNT_KHR, cl_uint, 1); + + // Confirm that querying CL_SEMAPHORE_PAYLOAD_KHR returns the unsignaled + // state + SEMAPHORE_PARAM_TEST(CL_SEMAPHORE_PAYLOAD_KHR, cl_semaphore_payload_khr, + 0); + } return TEST_PASS; } @@ -174,69 +193,83 @@ int test_external_semaphores_multi_context(cl_device_id deviceID, GET_PFN(deviceID, clEnqueueSignalSemaphoresKHR); GET_PFN(deviceID, clEnqueueWaitSemaphoresKHR); - const std::vector - vkExternalMemoryHandleTypeList = - getSupportedVulkanExternalMemoryHandleTypeList(); - VulkanExternalSemaphoreHandleType vkExternalSemaphoreHandleType = - getSupportedVulkanExternalSemaphoreHandleTypeList()[0]; - VulkanSemaphore vkVk2CLSemaphore(vkDevice, vkExternalSemaphoreHandleType); + std::vector + vkExternalSemaphoreHandleTypeList = + getSupportedInteropExternalSemaphoreHandleTypes(deviceID, vkDevice); - cl_int err = CL_SUCCESS; - - cl_context context2 = - clCreateContext(NULL, 1, &deviceID, notify_callback, NULL, &err); - if (!context2) + if (vkExternalSemaphoreHandleTypeList.empty()) { - print_error(err, "Unable to create testing context"); - return TEST_FAIL; + test_fail("No external semaphore handle types found\n"); } - clExternalSemaphore sema_ext_1(vkVk2CLSemaphore, context, - vkExternalSemaphoreHandleType, deviceID); - clExternalSemaphore sema_ext_2(vkVk2CLSemaphore, context2, - vkExternalSemaphoreHandleType, deviceID); - - clCommandQueueWrapper queue1 = - clCreateCommandQueue(context, deviceID, 0, &err); - test_error(err, "Could not create command queue"); - - clCommandQueueWrapper queue2 = - clCreateCommandQueue(context2, deviceID, 0, &err); - test_error(err, "Could not create command queue"); - - // Signal semaphore 1 and 2 - clEventWrapper signal_event; - err = clEnqueueSignalSemaphoresKHR(queue1, 1, &sema_ext_1.getCLSemaphore(), - nullptr, 0, nullptr, &signal_event); - test_error(err, "Could not signal semaphore"); + for (VulkanExternalSemaphoreHandleType vkExternalSemaphoreHandleType : + vkExternalSemaphoreHandleTypeList) + { + log_info_semaphore_type(vkExternalSemaphoreHandleType); + VulkanSemaphore vkVk2CLSemaphore(vkDevice, + vkExternalSemaphoreHandleType); + + cl_int err = CL_SUCCESS; + + cl_context context2 = + clCreateContext(NULL, 1, &deviceID, notify_callback, NULL, &err); + if (!context2) + { + print_error(err, "Unable to create testing context"); + return TEST_FAIL; + } + + clExternalSemaphore sema_ext_1(vkVk2CLSemaphore, context, + vkExternalSemaphoreHandleType, deviceID); + clExternalSemaphore sema_ext_2(vkVk2CLSemaphore, context2, + vkExternalSemaphoreHandleType, deviceID); + + clCommandQueueWrapper queue1 = + clCreateCommandQueue(context, deviceID, 0, &err); + test_error(err, "Could not create command queue"); + + clCommandQueueWrapper queue2 = + clCreateCommandQueue(context2, deviceID, 0, &err); + test_error(err, "Could not create command queue"); + + // Signal semaphore 1 and 2 + clEventWrapper signal_event; + err = clEnqueueSignalSemaphoresKHR(queue1, 1, + &sema_ext_1.getCLSemaphore(), + nullptr, 0, nullptr, &signal_event); + test_error(err, "Could not signal semaphore"); - // Wait semaphore 1 - clEventWrapper wait_1_event; - err = clEnqueueWaitSemaphoresKHR(queue1, 1, &sema_ext_1.getCLSemaphore(), - nullptr, 0, nullptr, &wait_1_event); - test_error(err, "Could not wait semaphore"); + // Wait semaphore 1 + clEventWrapper wait_1_event; + err = + clEnqueueWaitSemaphoresKHR(queue1, 1, &sema_ext_1.getCLSemaphore(), + nullptr, 0, nullptr, &wait_1_event); + test_error(err, "Could not wait semaphore"); - err = clEnqueueSignalSemaphoresKHR(queue2, 1, &sema_ext_2.getCLSemaphore(), - nullptr, 0, nullptr, &signal_event); - test_error(err, "Could not signal semaphore"); + err = clEnqueueSignalSemaphoresKHR(queue2, 1, + &sema_ext_2.getCLSemaphore(), + nullptr, 0, nullptr, &signal_event); + test_error(err, "Could not signal semaphore"); - // Wait semaphore 2 - clEventWrapper wait_2_event; - err = clEnqueueWaitSemaphoresKHR(queue2, 1, &sema_ext_2.getCLSemaphore(), - nullptr, 0, nullptr, &wait_2_event); - test_error(err, "Could not wait semaphore"); + // Wait semaphore 2 + clEventWrapper wait_2_event; + err = + clEnqueueWaitSemaphoresKHR(queue2, 1, &sema_ext_2.getCLSemaphore(), + nullptr, 0, nullptr, &wait_2_event); + test_error(err, "Could not wait semaphore"); - // Finish - err = clFinish(queue1); - test_error(err, "Could not finish queue"); + // Finish + err = clFinish(queue1); + test_error(err, "Could not finish queue"); - err = clFinish(queue2); - test_error(err, "Could not finish queue"); + err = clFinish(queue2); + test_error(err, "Could not finish queue"); - // Ensure all events are completed - test_assert_event_complete(signal_event); - test_assert_event_complete(wait_1_event); - test_assert_event_complete(wait_2_event); + // Ensure all events are completed + test_assert_event_complete(signal_event); + test_assert_event_complete(wait_1_event); + test_assert_event_complete(wait_2_event); + } return TEST_PASS; } @@ -264,44 +297,55 @@ static int semaphore_external_cross_queue_helper(cl_device_id deviceID, VulkanDevice vkDevice; - const std::vector - vkExternalMemoryHandleTypeList = - getSupportedVulkanExternalMemoryHandleTypeList(); - VulkanExternalSemaphoreHandleType vkExternalSemaphoreHandleType = - getSupportedVulkanExternalSemaphoreHandleTypeList()[0]; - VulkanSemaphore vkVk2CLSemaphore(vkDevice, vkExternalSemaphoreHandleType); + std::vector + vkExternalSemaphoreHandleTypeList = + getSupportedInteropExternalSemaphoreHandleTypes(deviceID, vkDevice); - clExternalSemaphore sema_ext(vkVk2CLSemaphore, context, - vkExternalSemaphoreHandleType, deviceID); + if (vkExternalSemaphoreHandleTypeList.empty()) + { + test_fail("No external semaphore handle types found\n"); + } - // Obtain pointers to semaphore's API - GET_PFN(deviceID, clEnqueueSignalSemaphoresKHR); - GET_PFN(deviceID, clEnqueueWaitSemaphoresKHR); + for (VulkanExternalSemaphoreHandleType vkExternalSemaphoreHandleType : + vkExternalSemaphoreHandleTypeList) + { + log_info_semaphore_type(vkExternalSemaphoreHandleType); + VulkanSemaphore vkVk2CLSemaphore(vkDevice, + vkExternalSemaphoreHandleType); - cl_int err = CL_SUCCESS; + clExternalSemaphore sema_ext(vkVk2CLSemaphore, context, + vkExternalSemaphoreHandleType, deviceID); - // Signal semaphore on queue_1 - clEventWrapper signal_event; - err = clEnqueueSignalSemaphoresKHR(queue_1, 1, &sema_ext.getCLSemaphore(), - nullptr, 0, nullptr, &signal_event); - test_error(err, "Could not signal semaphore"); + // Obtain pointers to semaphore's API + GET_PFN(deviceID, clEnqueueSignalSemaphoresKHR); + GET_PFN(deviceID, clEnqueueWaitSemaphoresKHR); - // Wait semaphore on queue_2 - clEventWrapper wait_event; - err = clEnqueueWaitSemaphoresKHR(queue_2, 1, &sema_ext.getCLSemaphore(), - nullptr, 0, nullptr, &wait_event); - test_error(err, "Could not wait semaphore"); + cl_int err = CL_SUCCESS; - // Finish queue_1 and queue_2 - err = clFinish(queue_1); - test_error(err, "Could not finish queue"); + // Signal semaphore on queue_1 + clEventWrapper signal_event; + err = + clEnqueueSignalSemaphoresKHR(queue_1, 1, &sema_ext.getCLSemaphore(), + nullptr, 0, nullptr, &signal_event); + test_error(err, "Could not signal semaphore"); + + // Wait semaphore on queue_2 + clEventWrapper wait_event; + err = clEnqueueWaitSemaphoresKHR(queue_2, 1, &sema_ext.getCLSemaphore(), + nullptr, 0, nullptr, &wait_event); + test_error(err, "Could not wait semaphore"); - err = clFinish(queue_2); - test_error(err, "Could not finish queue"); + // Finish queue_1 and queue_2 + err = clFinish(queue_1); + test_error(err, "Could not finish queue"); - // Ensure all events are completed - test_assert_event_complete(signal_event); - test_assert_event_complete(wait_event); + err = clFinish(queue_2); + test_error(err, "Could not finish queue"); + + // Ensure all events are completed + test_assert_event_complete(signal_event); + test_assert_event_complete(wait_event); + } return TEST_PASS; } @@ -331,42 +375,53 @@ int test_external_semaphores_simple_1(cl_device_id deviceID, cl_context context, GET_PFN(deviceID, clEnqueueSignalSemaphoresKHR); GET_PFN(deviceID, clEnqueueWaitSemaphoresKHR); - const std::vector - vkExternalMemoryHandleTypeList = - getSupportedVulkanExternalMemoryHandleTypeList(); - VulkanExternalSemaphoreHandleType vkExternalSemaphoreHandleType = - getSupportedVulkanExternalSemaphoreHandleTypeList()[0]; - VulkanSemaphore vkVk2CLSemaphore(vkDevice, vkExternalSemaphoreHandleType); + std::vector + vkExternalSemaphoreHandleTypeList = + getSupportedInteropExternalSemaphoreHandleTypes(deviceID, vkDevice); - clExternalSemaphore sema_ext(vkVk2CLSemaphore, context, - vkExternalSemaphoreHandleType, deviceID); + if (vkExternalSemaphoreHandleTypeList.empty()) + { + test_fail("No external semaphore handle types found\n"); + } - cl_int err = CL_SUCCESS; + for (VulkanExternalSemaphoreHandleType vkExternalSemaphoreHandleType : + vkExternalSemaphoreHandleTypeList) + { + log_info_semaphore_type(vkExternalSemaphoreHandleType); - // Create ooo queue - clCommandQueueWrapper queue = clCreateCommandQueue( - context, deviceID, CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE, &err); - test_error(err, "Could not create command queue"); + VulkanSemaphore vkVk2CLSemaphore(vkDevice, + vkExternalSemaphoreHandleType); + + clExternalSemaphore sema_ext(vkVk2CLSemaphore, context, + vkExternalSemaphoreHandleType, deviceID); + + cl_int err = CL_SUCCESS; - // Signal semaphore - clEventWrapper signal_event; - err = clEnqueueSignalSemaphoresKHR(queue, 1, &sema_ext.getCLSemaphore(), - nullptr, 0, nullptr, &signal_event); - test_error(err, "Could not signal semaphore"); + // Create ooo queue + clCommandQueueWrapper queue = clCreateCommandQueue( + context, deviceID, CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE, &err); + test_error(err, "Could not create command queue"); - // Wait semaphore - clEventWrapper wait_event; - err = clEnqueueWaitSemaphoresKHR(queue, 1, &sema_ext.getCLSemaphore(), - nullptr, 0, nullptr, &wait_event); - test_error(err, "Could not wait semaphore"); + // Signal semaphore + clEventWrapper signal_event; + err = clEnqueueSignalSemaphoresKHR(queue, 1, &sema_ext.getCLSemaphore(), + nullptr, 0, nullptr, &signal_event); + test_error(err, "Could not signal semaphore"); + + // Wait semaphore + clEventWrapper wait_event; + err = clEnqueueWaitSemaphoresKHR(queue, 1, &sema_ext.getCLSemaphore(), + nullptr, 0, nullptr, &wait_event); + test_error(err, "Could not wait semaphore"); - // Finish - err = clFinish(queue); - test_error(err, "Could not finish queue"); + // Finish + err = clFinish(queue); + test_error(err, "Could not finish queue"); - // Ensure all events are completed - test_assert_event_complete(signal_event); - test_assert_event_complete(wait_event); + // Ensure all events are completed + test_assert_event_complete(signal_event); + test_assert_event_complete(wait_event); + } return TEST_PASS; } @@ -397,73 +452,83 @@ int test_external_semaphores_simple_2(cl_device_id deviceID, cl_context context, GET_PFN(deviceID, clEnqueueSignalSemaphoresKHR); GET_PFN(deviceID, clEnqueueWaitSemaphoresKHR); - const std::vector - vkExternalMemoryHandleTypeList = - getSupportedVulkanExternalMemoryHandleTypeList(); - VulkanExternalSemaphoreHandleType vkExternalSemaphoreHandleType = - getSupportedVulkanExternalSemaphoreHandleTypeList()[0]; - VulkanSemaphore vkVk2CLSemaphore(vkDevice, vkExternalSemaphoreHandleType); + std::vector + vkExternalSemaphoreHandleTypeList = + getSupportedInteropExternalSemaphoreHandleTypes(deviceID, vkDevice); - clExternalSemaphore sema_ext(vkVk2CLSemaphore, context, - vkExternalSemaphoreHandleType, deviceID); + if (vkExternalSemaphoreHandleTypeList.empty()) + { + test_fail("No external semaphore handle types found\n"); + } - cl_int err = CL_SUCCESS; + for (VulkanExternalSemaphoreHandleType vkExternalSemaphoreHandleType : + vkExternalSemaphoreHandleTypeList) + { + log_info_semaphore_type(vkExternalSemaphoreHandleType); + VulkanSemaphore vkVk2CLSemaphore(vkDevice, + vkExternalSemaphoreHandleType); + + clExternalSemaphore sema_ext(vkVk2CLSemaphore, context, + vkExternalSemaphoreHandleType, deviceID); + + cl_int err = CL_SUCCESS; + + // Create ooo queue + clCommandQueueWrapper queue = clCreateCommandQueue( + context, deviceID, CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE, &err); + test_error(err, "Could not create command queue"); + + // Create user event + clEventWrapper user_event = clCreateUserEvent(context, &err); + test_error(err, "Could not create user event"); + + // Create Kernel + clProgramWrapper program; + clKernelWrapper kernel; + err = create_single_kernel_helper(context, &program, &kernel, 1, + &source, "empty"); + test_error(err, "Could not create kernel"); + + // Enqueue task_1 (dependency on user_event) + clEventWrapper task_1_event; + err = clEnqueueTask(queue, kernel, 1, &user_event, &task_1_event); + test_error(err, "Could not enqueue task 1"); + + // Signal semaphore + clEventWrapper signal_event; + err = clEnqueueSignalSemaphoresKHR(queue, 1, &sema_ext.getCLSemaphore(), + nullptr, 0, nullptr, &signal_event); + test_error(err, "Could not signal semaphore"); - // Create ooo queue - clCommandQueueWrapper queue = clCreateCommandQueue( - context, deviceID, CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE, &err); - test_error(err, "Could not create command queue"); + // Wait semaphore + clEventWrapper wait_event; + err = clEnqueueWaitSemaphoresKHR(queue, 1, &sema_ext.getCLSemaphore(), + nullptr, 0, nullptr, &wait_event); + test_error(err, "Could not wait semaphore"); + + // Flush and delay + err = clFlush(queue); + test_error(err, "Could not flush queue"); + std::this_thread::sleep_for(std::chrono::seconds(FLUSH_DELAY_S)); + + // Ensure all events are completed except for task_1 + test_assert_event_inprogress(task_1_event); + test_assert_event_complete(signal_event); + test_assert_event_complete(wait_event); + + // Complete user_event + err = clSetUserEventStatus(user_event, CL_COMPLETE); + test_error(err, "Could not set user event to CL_COMPLETE"); - // Create user event - clEventWrapper user_event = clCreateUserEvent(context, &err); - test_error(err, "Could not create user event"); - - // Create Kernel - clProgramWrapper program; - clKernelWrapper kernel; - err = create_single_kernel_helper(context, &program, &kernel, 1, &source, - "empty"); - test_error(err, "Could not create kernel"); - - // Enqueue task_1 (dependency on user_event) - clEventWrapper task_1_event; - err = clEnqueueTask(queue, kernel, 1, &user_event, &task_1_event); - test_error(err, "Could not enqueue task 1"); - - // Signal semaphore - clEventWrapper signal_event; - err = clEnqueueSignalSemaphoresKHR(queue, 1, &sema_ext.getCLSemaphore(), - nullptr, 0, nullptr, &signal_event); - test_error(err, "Could not signal semaphore"); - - // Wait semaphore - clEventWrapper wait_event; - err = clEnqueueWaitSemaphoresKHR(queue, 1, &sema_ext.getCLSemaphore(), - nullptr, 0, nullptr, &wait_event); - test_error(err, "Could not wait semaphore"); - - // Flush and delay - err = clFlush(queue); - test_error(err, "Could not flush queue"); - std::this_thread::sleep_for(std::chrono::seconds(FLUSH_DELAY_S)); - - // Ensure all events are completed except for task_1 - test_assert_event_inprogress(task_1_event); - test_assert_event_complete(signal_event); - test_assert_event_complete(wait_event); - - // Complete user_event - err = clSetUserEventStatus(user_event, CL_COMPLETE); - test_error(err, "Could not set user event to CL_COMPLETE"); - - // Finish - err = clFinish(queue); - test_error(err, "Could not finish queue"); - - // Ensure all events are completed - test_assert_event_complete(task_1_event); - test_assert_event_complete(signal_event); - test_assert_event_complete(wait_event); + // Finish + err = clFinish(queue); + test_error(err, "Could not finish queue"); + + // Ensure all events are completed + test_assert_event_complete(task_1_event); + test_assert_event_complete(signal_event); + test_assert_event_complete(wait_event); + } return TEST_PASS; } @@ -493,87 +558,97 @@ int test_external_semaphores_reuse(cl_device_id deviceID, cl_context context, GET_PFN(deviceID, clEnqueueSignalSemaphoresKHR); GET_PFN(deviceID, clEnqueueWaitSemaphoresKHR); - const std::vector - vkExternalMemoryHandleTypeList = - getSupportedVulkanExternalMemoryHandleTypeList(); - VulkanExternalSemaphoreHandleType vkExternalSemaphoreHandleType = - getSupportedVulkanExternalSemaphoreHandleTypeList()[0]; - VulkanSemaphore vkVk2CLSemaphore(vkDevice, vkExternalSemaphoreHandleType); - - clExternalSemaphore sema_ext(vkVk2CLSemaphore, context, - vkExternalSemaphoreHandleType, deviceID); - - cl_int err = CL_SUCCESS; + std::vector + vkExternalSemaphoreHandleTypeList = + getSupportedInteropExternalSemaphoreHandleTypes(deviceID, vkDevice); - // Create ooo queue - clCommandQueueWrapper queue = clCreateCommandQueue( - context, deviceID, CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE, &err); - test_error(err, "Could not create command queue"); + if (vkExternalSemaphoreHandleTypeList.empty()) + { + test_fail("No external semaphore handle types found\n"); + } - // Create Kernel - clProgramWrapper program; - clKernelWrapper kernel; - err = create_single_kernel_helper(context, &program, &kernel, 1, &source, - "empty"); - test_error(err, "Could not create kernel"); + for (VulkanExternalSemaphoreHandleType vkExternalSemaphoreHandleType : + vkExternalSemaphoreHandleTypeList) + { + log_info_semaphore_type(vkExternalSemaphoreHandleType); + VulkanSemaphore vkVk2CLSemaphore(vkDevice, + vkExternalSemaphoreHandleType); - constexpr size_t loop_count = 10; - clEventWrapper signal_events[loop_count]; - clEventWrapper wait_events[loop_count]; - clEventWrapper task_events[loop_count]; + clExternalSemaphore sema_ext(vkVk2CLSemaphore, context, + vkExternalSemaphoreHandleType, deviceID); - // Enqueue task_1 - err = clEnqueueTask(queue, kernel, 0, nullptr, &task_events[0]); - test_error(err, "Unable to enqueue task_1"); + cl_int err = CL_SUCCESS; - // Signal semaphore (dependency on task_1) - err = clEnqueueSignalSemaphoresKHR(queue, 1, &sema_ext.getCLSemaphore(), - nullptr, 1, &task_events[0], - &signal_events[0]); - test_error(err, "Could not signal semaphore"); + // Create ooo queue + clCommandQueueWrapper queue = clCreateCommandQueue( + context, deviceID, CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE, &err); + test_error(err, "Could not create command queue"); - // In a loop - size_t loop; - for (loop = 1; loop < loop_count; ++loop) - { - // Wait semaphore - err = clEnqueueWaitSemaphoresKHR(queue, 1, &sema_ext.getCLSemaphore(), - nullptr, 0, nullptr, - &wait_events[loop - 1]); - test_error(err, "Could not wait semaphore"); + // Create Kernel + clProgramWrapper program; + clKernelWrapper kernel; + err = create_single_kernel_helper(context, &program, &kernel, 1, + &source, "empty"); + test_error(err, "Could not create kernel"); - // Enqueue task_loop (dependency on wait) - err = clEnqueueTask(queue, kernel, 1, &wait_events[loop - 1], - &task_events[loop]); - test_error(err, "Unable to enqueue task_loop"); + constexpr size_t loop_count = 10; + clEventWrapper signal_events[loop_count]; + clEventWrapper wait_events[loop_count]; + clEventWrapper task_events[loop_count]; - // Wait for the "wait semaphore" to complete - err = clWaitForEvents(1, &wait_events[loop - 1]); - test_error(err, "Unable to wait for wait semaphore to complete"); + // Enqueue task_1 + err = clEnqueueTask(queue, kernel, 0, nullptr, &task_events[0]); + test_error(err, "Unable to enqueue task_1"); - // Signal semaphore (dependency on task_loop) + // Signal semaphore (dependency on task_1) err = clEnqueueSignalSemaphoresKHR(queue, 1, &sema_ext.getCLSemaphore(), - nullptr, 1, &task_events[loop], - &signal_events[loop]); + nullptr, 1, &task_events[0], + &signal_events[0]); test_error(err, "Could not signal semaphore"); - } - // Wait semaphore - err = - clEnqueueWaitSemaphoresKHR(queue, 1, &sema_ext.getCLSemaphore(), - nullptr, 0, nullptr, &wait_events[loop - 1]); - test_error(err, "Could not wait semaphore"); + // In a loop + size_t loop; + for (loop = 1; loop < loop_count; ++loop) + { + // Wait semaphore + err = clEnqueueWaitSemaphoresKHR( + queue, 1, &sema_ext.getCLSemaphore(), nullptr, 0, nullptr, + &wait_events[loop - 1]); + test_error(err, "Could not wait semaphore"); + + // Enqueue task_loop (dependency on wait) + err = clEnqueueTask(queue, kernel, 1, &wait_events[loop - 1], + &task_events[loop]); + test_error(err, "Unable to enqueue task_loop"); + + // Wait for the "wait semaphore" to complete + err = clWaitForEvents(1, &wait_events[loop - 1]); + test_error(err, "Unable to wait for wait semaphore to complete"); + + // Signal semaphore (dependency on task_loop) + err = clEnqueueSignalSemaphoresKHR( + queue, 1, &sema_ext.getCLSemaphore(), nullptr, 1, + &task_events[loop], &signal_events[loop]); + test_error(err, "Could not signal semaphore"); + } - // Finish - err = clFinish(queue); - test_error(err, "Could not finish queue"); + // Wait semaphore + err = clEnqueueWaitSemaphoresKHR(queue, 1, &sema_ext.getCLSemaphore(), + nullptr, 0, nullptr, + &wait_events[loop - 1]); + test_error(err, "Could not wait semaphore"); - // Ensure all events are completed - for (loop = 0; loop < loop_count; ++loop) - { - test_assert_event_complete(wait_events[loop]); - test_assert_event_complete(signal_events[loop]); - test_assert_event_complete(task_events[loop]); + // Finish + err = clFinish(queue); + test_error(err, "Could not finish queue"); + + // Ensure all events are completed + for (loop = 0; loop < loop_count; ++loop) + { + test_assert_event_complete(wait_events[loop]); + test_assert_event_complete(signal_events[loop]); + test_assert_event_complete(task_events[loop]); + } } return TEST_PASS; @@ -606,40 +681,51 @@ static int external_semaphore_cross_queue_helper(cl_device_id deviceID, GET_PFN(deviceID, clEnqueueSignalSemaphoresKHR); GET_PFN(deviceID, clEnqueueWaitSemaphoresKHR); - const std::vector - vkExternalMemoryHandleTypeList = - getSupportedVulkanExternalMemoryHandleTypeList(); - VulkanExternalSemaphoreHandleType vkExternalSemaphoreHandleType = - getSupportedVulkanExternalSemaphoreHandleTypeList()[0]; - VulkanSemaphore vkVk2CLSemaphore(vkDevice, vkExternalSemaphoreHandleType); + std::vector + vkExternalSemaphoreHandleTypeList = + getSupportedInteropExternalSemaphoreHandleTypes(deviceID, vkDevice); - clExternalSemaphore sema_ext(vkVk2CLSemaphore, context, - vkExternalSemaphoreHandleType, deviceID); + if (vkExternalSemaphoreHandleTypeList.empty()) + { + test_fail("No external semaphore handle types found\n"); + } - cl_int err = CL_SUCCESS; + for (VulkanExternalSemaphoreHandleType vkExternalSemaphoreHandleType : + vkExternalSemaphoreHandleTypeList) + { + log_info_semaphore_type(vkExternalSemaphoreHandleType); + VulkanSemaphore vkVk2CLSemaphore(vkDevice, + vkExternalSemaphoreHandleType); - // Signal semaphore on queue_1 - clEventWrapper signal_event; - err = clEnqueueSignalSemaphoresKHR(queue_1, 1, &sema_ext.getCLSemaphore(), - nullptr, 0, nullptr, &signal_event); - test_error(err, "Could not signal semaphore"); + clExternalSemaphore sema_ext(vkVk2CLSemaphore, context, + vkExternalSemaphoreHandleType, deviceID); - // Wait semaphore on queue_2 - clEventWrapper wait_event; - err = clEnqueueWaitSemaphoresKHR(queue_2, 1, &sema_ext.getCLSemaphore(), - nullptr, 0, nullptr, &wait_event); - test_error(err, "Could not wait semaphore"); + cl_int err = CL_SUCCESS; - // Finish queue_1 and queue_2 - err = clFinish(queue_1); - test_error(err, "Could not finish queue"); + // Signal semaphore on queue_1 + clEventWrapper signal_event; + err = + clEnqueueSignalSemaphoresKHR(queue_1, 1, &sema_ext.getCLSemaphore(), + nullptr, 0, nullptr, &signal_event); + test_error(err, "Could not signal semaphore"); + + // Wait semaphore on queue_2 + clEventWrapper wait_event; + err = clEnqueueWaitSemaphoresKHR(queue_2, 1, &sema_ext.getCLSemaphore(), + nullptr, 0, nullptr, &wait_event); + test_error(err, "Could not wait semaphore"); + + // Finish queue_1 and queue_2 + err = clFinish(queue_1); + test_error(err, "Could not finish queue"); - err = clFinish(queue_2); - test_error(err, "Could not finish queue"); + err = clFinish(queue_2); + test_error(err, "Could not finish queue"); - // Ensure all events are completed - test_assert_event_complete(signal_event); - test_assert_event_complete(wait_event); + // Ensure all events are completed + test_assert_event_complete(signal_event); + test_assert_event_complete(wait_event); + } return TEST_PASS; } @@ -721,59 +807,73 @@ int test_external_semaphores_cross_queues_io2(cl_device_id deviceID, GET_PFN(deviceID, clEnqueueSignalSemaphoresKHR); GET_PFN(deviceID, clEnqueueWaitSemaphoresKHR); - const std::vector - vkExternalMemoryHandleTypeList = - getSupportedVulkanExternalMemoryHandleTypeList(); - VulkanExternalSemaphoreHandleType vkExternalSemaphoreHandleType = - getSupportedVulkanExternalSemaphoreHandleTypeList()[0]; - VulkanSemaphore vkVk2CLSemaphore(vkDevice, vkExternalSemaphoreHandleType); - - clExternalSemaphore sema_ext_1(vkVk2CLSemaphore, context, - vkExternalSemaphoreHandleType, deviceID); - clExternalSemaphore sema_ext_2(vkVk2CLSemaphore, context2, - vkExternalSemaphoreHandleType, deviceID); + std::vector + vkExternalSemaphoreHandleTypeList = + getSupportedInteropExternalSemaphoreHandleTypes(deviceID, vkDevice); - clCommandQueueWrapper queue1 = - clCreateCommandQueue(context, deviceID, 0, &err); - test_error(err, "Could not create command queue"); - - clCommandQueueWrapper queue2 = - clCreateCommandQueue(context2, deviceID, 0, &err); - test_error(err, "Could not create command queue"); + if (vkExternalSemaphoreHandleTypeList.empty()) + { + test_fail("No external semaphore handle types found\n"); + } - // Signal semaphore 1 and 2 - clEventWrapper signal_event; - err = clEnqueueSignalSemaphoresKHR(queue1, 1, &sema_ext_1.getCLSemaphore(), - nullptr, 0, nullptr, &signal_event); - test_error(err, "Could not signal semaphore"); + for (VulkanExternalSemaphoreHandleType vkExternalSemaphoreHandleType : + vkExternalSemaphoreHandleTypeList) + { + log_info_semaphore_type(vkExternalSemaphoreHandleType); + VulkanSemaphore vkVk2CLSemaphore(vkDevice, + vkExternalSemaphoreHandleType); + + clExternalSemaphore sema_ext_1(vkVk2CLSemaphore, context, + vkExternalSemaphoreHandleType, deviceID); + clExternalSemaphore sema_ext_2(vkVk2CLSemaphore, context2, + vkExternalSemaphoreHandleType, deviceID); + + clCommandQueueWrapper queue1 = + clCreateCommandQueue(context, deviceID, 0, &err); + test_error(err, "Could not create command queue"); + + clCommandQueueWrapper queue2 = + clCreateCommandQueue(context2, deviceID, 0, &err); + test_error(err, "Could not create command queue"); + + // Signal semaphore 1 and 2 + clEventWrapper signal_event; + err = clEnqueueSignalSemaphoresKHR(queue1, 1, + &sema_ext_1.getCLSemaphore(), + nullptr, 0, nullptr, &signal_event); + test_error(err, "Could not signal semaphore"); - // Wait semaphore 1 - clEventWrapper wait_1_event; - err = clEnqueueWaitSemaphoresKHR(queue1, 1, &sema_ext_1.getCLSemaphore(), - nullptr, 0, nullptr, &wait_1_event); - test_error(err, "Could not wait semaphore"); + // Wait semaphore 1 + clEventWrapper wait_1_event; + err = + clEnqueueWaitSemaphoresKHR(queue1, 1, &sema_ext_1.getCLSemaphore(), + nullptr, 0, nullptr, &wait_1_event); + test_error(err, "Could not wait semaphore"); - err = clEnqueueSignalSemaphoresKHR(queue2, 1, &sema_ext_2.getCLSemaphore(), - nullptr, 0, nullptr, &signal_event); - test_error(err, "Could not signal semaphore"); + err = clEnqueueSignalSemaphoresKHR(queue2, 1, + &sema_ext_2.getCLSemaphore(), + nullptr, 0, nullptr, &signal_event); + test_error(err, "Could not signal semaphore"); - // Wait semaphore 2 - clEventWrapper wait_2_event; - err = clEnqueueWaitSemaphoresKHR(queue2, 1, &sema_ext_2.getCLSemaphore(), - nullptr, 0, nullptr, &wait_2_event); - test_error(err, "Could not wait semaphore"); + // Wait semaphore 2 + clEventWrapper wait_2_event; + err = + clEnqueueWaitSemaphoresKHR(queue2, 1, &sema_ext_2.getCLSemaphore(), + nullptr, 0, nullptr, &wait_2_event); + test_error(err, "Could not wait semaphore"); - // Finish - err = clFinish(queue1); - test_error(err, "Could not finish queue"); + // Finish + err = clFinish(queue1); + test_error(err, "Could not finish queue"); - err = clFinish(queue2); - test_error(err, "Could not finish queue"); + err = clFinish(queue2); + test_error(err, "Could not finish queue"); - // Ensure all events are completed - test_assert_event_complete(signal_event); - test_assert_event_complete(wait_1_event); - test_assert_event_complete(wait_2_event); + // Ensure all events are completed + test_assert_event_complete(signal_event); + test_assert_event_complete(wait_1_event); + test_assert_event_complete(wait_2_event); + } return TEST_PASS; } @@ -804,54 +904,65 @@ int test_external_semaphores_multi_signal(cl_device_id deviceID, GET_PFN(deviceID, clEnqueueSignalSemaphoresKHR); GET_PFN(deviceID, clEnqueueWaitSemaphoresKHR); - const std::vector - vkExternalMemoryHandleTypeList = - getSupportedVulkanExternalMemoryHandleTypeList(); - VulkanExternalSemaphoreHandleType vkExternalSemaphoreHandleType = - getSupportedVulkanExternalSemaphoreHandleTypeList()[0]; - VulkanSemaphore vkVk2CLSemaphore1(vkDevice, vkExternalSemaphoreHandleType); - VulkanSemaphore vkVk2CLSemaphore2(vkDevice, vkExternalSemaphoreHandleType); + std::vector + vkExternalSemaphoreHandleTypeList = + getSupportedInteropExternalSemaphoreHandleTypes(deviceID, vkDevice); - clExternalSemaphore sema_ext_1(vkVk2CLSemaphore1, context, - vkExternalSemaphoreHandleType, deviceID); - clExternalSemaphore sema_ext_2(vkVk2CLSemaphore2, context, - vkExternalSemaphoreHandleType, deviceID); + if (vkExternalSemaphoreHandleTypeList.empty()) + { + test_fail("No external semaphore handle types found\n"); + } - cl_int err = CL_SUCCESS; + for (VulkanExternalSemaphoreHandleType vkExternalSemaphoreHandleType : + vkExternalSemaphoreHandleTypeList) + { + log_info_semaphore_type(vkExternalSemaphoreHandleType); + VulkanSemaphore vkVk2CLSemaphore1(vkDevice, + vkExternalSemaphoreHandleType); + VulkanSemaphore vkVk2CLSemaphore2(vkDevice, + vkExternalSemaphoreHandleType); + + clExternalSemaphore sema_ext_1(vkVk2CLSemaphore1, context, + vkExternalSemaphoreHandleType, deviceID); + clExternalSemaphore sema_ext_2(vkVk2CLSemaphore2, context, + vkExternalSemaphoreHandleType, deviceID); + + cl_int err = CL_SUCCESS; + + // Create ooo queue + clCommandQueueWrapper queue = clCreateCommandQueue( + context, deviceID, CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE, &err); + test_error(err, "Could not create command queue"); + + // Signal semaphore 1 and 2 + clEventWrapper signal_event; + cl_semaphore_khr sema_list[] = { sema_ext_1.getCLSemaphore(), + sema_ext_2.getCLSemaphore() }; + err = clEnqueueSignalSemaphoresKHR(queue, 2, sema_list, nullptr, 0, + nullptr, &signal_event); + test_error(err, "Could not signal semaphore"); - // Create ooo queue - clCommandQueueWrapper queue = clCreateCommandQueue( - context, deviceID, CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE, &err); - test_error(err, "Could not create command queue"); + // Wait semaphore 1 + clEventWrapper wait_1_event; + err = clEnqueueWaitSemaphoresKHR(queue, 1, &sema_ext_1.getCLSemaphore(), + nullptr, 0, nullptr, &wait_1_event); + test_error(err, "Could not wait semaphore"); + + // Wait semaphore 2 + clEventWrapper wait_2_event; + err = clEnqueueWaitSemaphoresKHR(queue, 1, &sema_ext_2.getCLSemaphore(), + nullptr, 0, nullptr, &wait_2_event); + test_error(err, "Could not wait semaphore"); - // Signal semaphore 1 and 2 - clEventWrapper signal_event; - cl_semaphore_khr sema_list[] = { sema_ext_1.getCLSemaphore(), - sema_ext_2.getCLSemaphore() }; - err = clEnqueueSignalSemaphoresKHR(queue, 2, sema_list, nullptr, 0, nullptr, - &signal_event); - test_error(err, "Could not signal semaphore"); - - // Wait semaphore 1 - clEventWrapper wait_1_event; - err = clEnqueueWaitSemaphoresKHR(queue, 1, &sema_ext_1.getCLSemaphore(), - nullptr, 0, nullptr, &wait_1_event); - test_error(err, "Could not wait semaphore"); - - // Wait semaphore 2 - clEventWrapper wait_2_event; - err = clEnqueueWaitSemaphoresKHR(queue, 1, &sema_ext_2.getCLSemaphore(), - nullptr, 0, nullptr, &wait_2_event); - test_error(err, "Could not wait semaphore"); - - // Finish - err = clFinish(queue); - test_error(err, "Could not finish queue"); - - // Ensure all events are completed - test_assert_event_complete(signal_event); - test_assert_event_complete(wait_1_event); - test_assert_event_complete(wait_2_event); + // Finish + err = clFinish(queue); + test_error(err, "Could not finish queue"); + + // Ensure all events are completed + test_assert_event_complete(signal_event); + test_assert_event_complete(wait_1_event); + test_assert_event_complete(wait_2_event); + } return TEST_PASS; } @@ -882,54 +993,67 @@ int test_external_semaphores_multi_wait(cl_device_id deviceID, GET_PFN(deviceID, clEnqueueSignalSemaphoresKHR); GET_PFN(deviceID, clEnqueueWaitSemaphoresKHR); - const std::vector - vkExternalMemoryHandleTypeList = - getSupportedVulkanExternalMemoryHandleTypeList(); - VulkanExternalSemaphoreHandleType vkExternalSemaphoreHandleType = - getSupportedVulkanExternalSemaphoreHandleTypeList()[0]; - VulkanSemaphore vkVk2CLSemaphore1(vkDevice, vkExternalSemaphoreHandleType); - VulkanSemaphore vkVk2CLSemaphore2(vkDevice, vkExternalSemaphoreHandleType); + std::vector + vkExternalSemaphoreHandleTypeList = + getSupportedInteropExternalSemaphoreHandleTypes(deviceID, vkDevice); - clExternalSemaphore sema_ext_1(vkVk2CLSemaphore1, context, - vkExternalSemaphoreHandleType, deviceID); - clExternalSemaphore sema_ext_2(vkVk2CLSemaphore2, context, - vkExternalSemaphoreHandleType, deviceID); + if (vkExternalSemaphoreHandleTypeList.empty()) + { + test_fail("No external semaphore handle types found\n"); + } - cl_int err = CL_SUCCESS; + for (VulkanExternalSemaphoreHandleType vkExternalSemaphoreHandleType : + vkExternalSemaphoreHandleTypeList) + { + log_info_semaphore_type(vkExternalSemaphoreHandleType); + VulkanSemaphore vkVk2CLSemaphore1(vkDevice, + vkExternalSemaphoreHandleType); + VulkanSemaphore vkVk2CLSemaphore2(vkDevice, + vkExternalSemaphoreHandleType); + + clExternalSemaphore sema_ext_1(vkVk2CLSemaphore1, context, + vkExternalSemaphoreHandleType, deviceID); + clExternalSemaphore sema_ext_2(vkVk2CLSemaphore2, context, + vkExternalSemaphoreHandleType, deviceID); + + cl_int err = CL_SUCCESS; + + // Create ooo queue + clCommandQueueWrapper queue = clCreateCommandQueue( + context, deviceID, CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE, &err); + test_error(err, "Could not create command queue"); + + // Signal semaphore 1 + clEventWrapper signal_1_event; + err = + clEnqueueSignalSemaphoresKHR(queue, 1, &sema_ext_1.getCLSemaphore(), + nullptr, 0, nullptr, &signal_1_event); + test_error(err, "Could not signal semaphore"); - // Create ooo queue - clCommandQueueWrapper queue = clCreateCommandQueue( - context, deviceID, CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE, &err); - test_error(err, "Could not create command queue"); + // Signal semaphore 2 + clEventWrapper signal_2_event; + err = + clEnqueueSignalSemaphoresKHR(queue, 1, &sema_ext_2.getCLSemaphore(), + nullptr, 0, nullptr, &signal_2_event); + test_error(err, "Could not signal semaphore"); - // Signal semaphore 1 - clEventWrapper signal_1_event; - err = clEnqueueSignalSemaphoresKHR(queue, 1, &sema_ext_1.getCLSemaphore(), - nullptr, 0, nullptr, &signal_1_event); - test_error(err, "Could not signal semaphore"); - - // Signal semaphore 2 - clEventWrapper signal_2_event; - err = clEnqueueSignalSemaphoresKHR(queue, 1, &sema_ext_2.getCLSemaphore(), - nullptr, 0, nullptr, &signal_2_event); - test_error(err, "Could not signal semaphore"); - - // Wait semaphore 1 and 2 - clEventWrapper wait_event; - cl_semaphore_khr sema_list[] = { sema_ext_1.getCLSemaphore(), - sema_ext_2.getCLSemaphore() }; - err = clEnqueueWaitSemaphoresKHR(queue, 2, sema_list, nullptr, 0, nullptr, - &wait_event); - test_error(err, "Could not wait semaphore"); - - // Finish - err = clFinish(queue); - test_error(err, "Could not finish queue"); - - // Ensure all events are completed - test_assert_event_complete(signal_1_event); - test_assert_event_complete(signal_2_event); - test_assert_event_complete(wait_event); + // Wait semaphore 1 and 2 + clEventWrapper wait_event; + cl_semaphore_khr sema_list[] = { sema_ext_1.getCLSemaphore(), + sema_ext_2.getCLSemaphore() }; + err = clEnqueueWaitSemaphoresKHR(queue, 2, sema_list, nullptr, 0, + nullptr, &wait_event); + test_error(err, "Could not wait semaphore"); + + // Finish + err = clFinish(queue); + test_error(err, "Could not finish queue"); + + // Ensure all events are completed + test_assert_event_complete(signal_1_event); + test_assert_event_complete(signal_2_event); + test_assert_event_complete(wait_event); + } return TEST_PASS; } diff --git a/test_conformance/images/kernel_read_write/test_cl_ext_image_from_buffer.cpp b/test_conformance/images/kernel_read_write/test_cl_ext_image_from_buffer.cpp index 6d94c2ff28..7b4860dbcb 100644 --- a/test_conformance/images/kernel_read_write/test_cl_ext_image_from_buffer.cpp +++ b/test_conformance/images/kernel_read_write/test_cl_ext_image_from_buffer.cpp @@ -18,6 +18,8 @@ #include "../common.h" #include "test_cl_ext_image_buffer.hpp" +static inline bool is_power_of_two(size_t num) { return !(num & (num - 1)); } + static int get_image_requirement_alignment( cl_device_id device, cl_context context, cl_mem_flags flags, const cl_image_format* image_format, const cl_image_desc* image_desc, @@ -79,11 +81,17 @@ int image2d_from_buffer_positive(cl_device_id device, cl_context context, return TEST_SKIPPED_ITSELF; } - std::vector imageTypes{ - CL_MEM_OBJECT_IMAGE1D, CL_MEM_OBJECT_IMAGE2D, - CL_MEM_OBJECT_IMAGE3D, CL_MEM_OBJECT_IMAGE1D_BUFFER, - CL_MEM_OBJECT_IMAGE1D_ARRAY, CL_MEM_OBJECT_IMAGE2D_ARRAY - }; + cl_uint row_pitch_alignment_2d = 0; + cl_int err = clGetDeviceInfo(device, CL_DEVICE_IMAGE_PITCH_ALIGNMENT, + sizeof(row_pitch_alignment_2d), + &row_pitch_alignment_2d, nullptr); + test_error(err, "Error clGetDeviceInfo"); + + cl_uint base_address_alignment_2d = 0; + err = clGetDeviceInfo(device, CL_DEVICE_IMAGE_BASE_ADDRESS_ALIGNMENT, + sizeof(base_address_alignment_2d), + &base_address_alignment_2d, nullptr); + test_error(err, "Error clGetDeviceInfo"); std::vector flagTypes{ CL_MEM_READ_ONLY, CL_MEM_WRITE_ONLY, CL_MEM_READ_WRITE, @@ -91,55 +99,43 @@ int image2d_from_buffer_positive(cl_device_id device, cl_context context, for (auto flagType : flagTypes) { - for (auto imageType : imageTypes) - { - /* Get the list of supported image formats */ - std::vector formatList; - if (TEST_PASS - != get_format_list(context, imageType, formatList, flagType) - || formatList.size() == 0) - { - test_fail("Failure to get supported formats list\n"); - } - - cl_uint row_pitch_alignment_2d = 0; - cl_int err = - clGetDeviceInfo(device, CL_DEVICE_IMAGE_PITCH_ALIGNMENT, - sizeof(row_pitch_alignment_2d), - &row_pitch_alignment_2d, nullptr); - test_error(err, "Error clGetDeviceInfo"); - cl_uint base_address_alignment_2d = 0; - err = - clGetDeviceInfo(device, CL_DEVICE_IMAGE_BASE_ADDRESS_ALIGNMENT, - sizeof(base_address_alignment_2d), - &base_address_alignment_2d, nullptr); - test_error(err, "Error clGetDeviceInfo"); + /* Get the list of supported image formats */ + std::vector formatList; + if (TEST_PASS + != get_format_list(context, CL_MEM_OBJECT_IMAGE2D, formatList, + flagType) + || formatList.size() == 0) + { + test_fail("Failure to get supported formats list\n"); + } - for (auto format : formatList) - { - cl_image_desc image_desc = { 0 }; - image_desc_init(&image_desc, imageType); + for (auto format : formatList) + { + cl_image_desc image_desc = { 0 }; + image_desc_init(&image_desc, CL_MEM_OBJECT_IMAGE2D); - cl_mem_flags flag = (flagType == CL_MEM_KERNEL_READ_AND_WRITE) - ? CL_MEM_READ_WRITE - : flagType; + cl_mem_flags flag = (flagType == CL_MEM_KERNEL_READ_AND_WRITE) + ? CL_MEM_READ_WRITE + : flagType; - size_t row_pitch_alignment = 0; - size_t base_address_alignment = 0; + size_t row_pitch_alignment = 0; + size_t base_address_alignment = 0; - int get_error = get_image_requirement_alignment( - device, context, flag, &format, &image_desc, - &row_pitch_alignment, nullptr, &base_address_alignment); - if (TEST_PASS != get_error) - { - return get_error; - } + int get_error = get_image_requirement_alignment( + device, context, flag, &format, &image_desc, + &row_pitch_alignment, nullptr, &base_address_alignment); + if (TEST_PASS != get_error) + { + return get_error; + } - const size_t element_size = - get_format_size(context, &format, imageType, flag); + const size_t element_size = + get_format_size(context, &format, CL_MEM_OBJECT_IMAGE2D, flag); - /* Alignements in pixels vs bytes */ + if (is_power_of_two(element_size)) + { + /* Alignments in pixels vs bytes */ if (base_address_alignment > base_address_alignment_2d * element_size) { diff --git a/test_conformance/profiling/main.cpp b/test_conformance/profiling/main.cpp index 6e59f611df..0498472836 100644 --- a/test_conformance/profiling/main.cpp +++ b/test_conformance/profiling/main.cpp @@ -17,6 +17,7 @@ #include #include +#include #include "procs.h" #include "harness/testHarness.h" @@ -72,12 +73,17 @@ int check_times(cl_ulong queueStart, cl_ulong commandSubmit, cl_ulong commandSta err = clGetDeviceInfo(device, CL_DEVICE_PROFILING_TIMER_RESOLUTION, sizeof(profiling_resolution), &profiling_resolution, NULL); test_error(err, "clGetDeviceInfo for CL_DEVICE_PROFILING_TIMER_RESOLUTION failed.\n"); - log_info("CL_PROFILING_COMMAND_QUEUED: %llu CL_PROFILING_COMMAND_SUBMIT: %llu CL_PROFILING_COMMAND_START: %llu CL_PROFILING_COMMAND_END: %llu CL_DEVICE_PROFILING_TIMER_RESOLUTION: %ld\n", - queueStart, commandSubmit, commandStart, commandEnd, profiling_resolution); - - double queueTosubmitTimeS = (double)(commandSubmit - queueStart)*1e-9; - double submitToStartTimeS = (double)(commandStart - commandSubmit)*1e-9; - double startToEndTimeS = (double)(commandEnd - commandStart)*1e-9; + log_info("CL_PROFILING_COMMAND_QUEUED: %" PRIu64 + " CL_PROFILING_COMMAND_SUBMIT: %" PRIu64 + " CL_PROFILING_COMMAND_START: %" PRIu64 + " CL_PROFILING_COMMAND_END: %" PRIu64 + " CL_DEVICE_PROFILING_TIMER_RESOLUTION: %zu\n", + queueStart, commandSubmit, commandStart, commandEnd, + profiling_resolution); + + double queueTosubmitTimeS = (double)(commandSubmit - queueStart) * 1e-9; + double submitToStartTimeS = (double)(commandStart - commandSubmit) * 1e-9; + double startToEndTimeS = (double)(commandEnd - commandStart) * 1e-9; log_info( "Profiling info:\n" ); log_info( "Time from queue to submit : %fms\n", (double)(queueTosubmitTimeS) * 1000.f ); diff --git a/test_conformance/subgroups/test_subgroup.cpp b/test_conformance/subgroups/test_subgroup.cpp index 75e9d4aeec..3b72913ee5 100644 --- a/test_conformance/subgroups/test_subgroup.cpp +++ b/test_conformance/subgroups/test_subgroup.cpp @@ -169,6 +169,7 @@ int test_subgroup_functions(cl_device_id device, cl_context context, constexpr size_t global_work_size = 2000; constexpr size_t local_work_size = 200; WorkGroupParams test_params(global_work_size, local_work_size); + test_params.use_core_subgroups = useCoreSubgroups; test_params.save_kernel_source(sub_group_reduction_scan_source); test_params.save_kernel_source(sub_group_generic_source, "sub_group_broadcast"); diff --git a/test_conformance/vulkan/main.cpp b/test_conformance/vulkan/main.cpp index eb1afeb0f6..5c699b6303 100644 --- a/test_conformance/vulkan/main.cpp +++ b/test_conformance/vulkan/main.cpp @@ -53,7 +53,7 @@ static void params_reset() extern int test_buffer_common(cl_device_id device_, cl_context context_, cl_command_queue queue_, int numElements_, - float use_fence); + bool use_fence); extern int test_image_common(cl_device_id device_, cl_context context_, cl_command_queue queue_, int numElements_); @@ -227,10 +227,6 @@ size_t parseParams(int argc, const char *argv[], const char **argList) { disableNTHandleType = true; } - if (!strcmp(argv[i], "--enableOffset")) - { - enableOffset = true; - } if (strcmp(argv[i], "-h") == 0) { printUsage(argv[0]); diff --git a/test_conformance/vulkan/shaders/buffer.comp b/test_conformance/vulkan/shaders/buffer.comp index d8756f9249..3e4eae55f9 100644 --- a/test_conformance/vulkan/shaders/buffer.comp +++ b/test_conformance/vulkan/shaders/buffer.comp @@ -15,7 +15,7 @@ layout(binding = 1) buffer Buffer { uint8_t ptr[]; } bufferPtrList[MAX_BUFFERS]; -layout(local_size_x = 512) in; +layout(local_size_x = 128) in; void main() { for (uint32_t bufIdx = 0; bufIdx < numBuffers; bufIdx++) { uint32_t ptrIdx = gl_GlobalInvocationID.x; diff --git a/test_conformance/vulkan/shaders/buffer.spv b/test_conformance/vulkan/shaders/buffer.spv index 685523ba5f..c9d15950c2 100644 Binary files a/test_conformance/vulkan/shaders/buffer.spv and b/test_conformance/vulkan/shaders/buffer.spv differ diff --git a/test_conformance/vulkan/test_vulkan_api_consistency.cpp b/test_conformance/vulkan/test_vulkan_api_consistency.cpp index d12b3bfe48..fe06052e10 100644 --- a/test_conformance/vulkan/test_vulkan_api_consistency.cpp +++ b/test_conformance/vulkan/test_vulkan_api_consistency.cpp @@ -93,9 +93,9 @@ int test_consistency_external_buffer(cl_device_id deviceID, cl_context _context, int fd; std::vector extMemProperties{ - (cl_mem_properties)CL_DEVICE_HANDLE_LIST_KHR, + (cl_mem_properties)CL_MEM_DEVICE_HANDLE_LIST_KHR, (cl_mem_properties)devList[0], - (cl_mem_properties)CL_DEVICE_HANDLE_LIST_END_KHR, + (cl_mem_properties)CL_MEM_DEVICE_HANDLE_LIST_END_KHR, }; cl_external_memory_handle_type_khr type; switch (vkExternalMemoryHandleType) @@ -162,9 +162,9 @@ int test_consistency_external_buffer(cl_device_id deviceID, cl_context _context, (cl_mem_properties)type, (cl_mem_properties)-64, // Passing random invalid fd #endif - (cl_mem_properties)CL_DEVICE_HANDLE_LIST_KHR, + (cl_mem_properties)CL_MEM_DEVICE_HANDLE_LIST_KHR, (cl_mem_properties)devList[0], - (cl_mem_properties)CL_DEVICE_HANDLE_LIST_END_KHR, + (cl_mem_properties)CL_MEM_DEVICE_HANDLE_LIST_END_KHR, 0 }; buffer = clCreateBufferWithProperties(context, extMemProperties2.data(), 1, @@ -219,9 +219,8 @@ int test_consistency_external_image(cl_device_id deviceID, cl_context _context, #else if (!is_extension_available(devList[0], "cl_khr_external_memory_opaque_fd")) { - throw std::runtime_error( - "Device does not support cl_khr_external_memory_opaque_fd " - "extension \n"); + test_fail("Device does not support cl_khr_external_memory_opaque_fd " + "extension \n"); } #endif uint32_t width = 256; @@ -257,9 +256,9 @@ int test_consistency_external_image(cl_device_id deviceID, cl_context _context, void* handle = NULL; int fd; std::vector extMemProperties{ - (cl_mem_properties)CL_DEVICE_HANDLE_LIST_KHR, + (cl_mem_properties)CL_MEM_DEVICE_HANDLE_LIST_KHR, (cl_mem_properties)devList[0], - (cl_mem_properties)CL_DEVICE_HANDLE_LIST_END_KHR, + (cl_mem_properties)CL_MEM_DEVICE_HANDLE_LIST_END_KHR, }; switch (vkExternalMemoryHandleType) { @@ -324,26 +323,6 @@ int test_consistency_external_image(cl_device_id deviceID, cl_context _context, test_error(errNum, "Unable to create Image with Properties"); image.reset(); - // Passing properties, image_desc and image_format all as NULL - image = clCreateImageWithProperties(context, NULL, CL_MEM_READ_WRITE, NULL, - NULL, NULL, &errNum); - test_failure_error( - errNum, CL_INVALID_IMAGE_DESCRIPTOR, - "Image creation must fail with CL_INVALID_IMAGE_DESCRIPTOR " - "when all are passed as NULL"); - - image.reset(); - - // Passing NULL properties and a valid image_format and image_desc - image = - clCreateImageWithProperties(context, NULL, CL_MEM_READ_WRITE, - &img_format, &image_desc, NULL, &errNum); - test_error(errNum, - "Unable to create image with NULL properties " - "with valid image format and image desc"); - - image.reset(); - // Passing image_format as NULL image = clCreateImageWithProperties(context, extMemProperties.data(), CL_MEM_READ_WRITE, NULL, &image_desc, @@ -396,103 +375,107 @@ int test_consistency_external_semaphore(cl_device_id deviceID, cl_device_id devList[] = { deviceID, NULL }; -#ifdef _WIN32 - if (!is_extension_available(devList[0], "cl_khr_external_semaphore_win32")) - { - throw std::runtime_error( - "Device does not support cl_khr_external_semaphore_win32 " - "extension \n"); - } -#else - if (!is_extension_available(devList[0], - "cl_khr_external_semaphore_opaque_fd")) + std::vector supportedExternalSemaphores = + getSupportedInteropExternalSemaphoreHandleTypes(devList[0], vkDevice); + + if (supportedExternalSemaphores.empty()) { - throw std::runtime_error( - "Device does not support " - "cl_khr_external_semaphore_opaque_fd extension \n"); + test_fail("No supported external semaphore types found\n"); } -#endif - VulkanExternalSemaphoreHandleType vkExternalSemaphoreHandleType = - getSupportedVulkanExternalSemaphoreHandleTypeList()[0]; - VulkanSemaphore vkVk2Clsemaphore(vkDevice, vkExternalSemaphoreHandleType); - VulkanSemaphore vkCl2Vksemaphore(vkDevice, vkExternalSemaphoreHandleType); - cl_semaphore_khr clCl2Vksemaphore; - cl_semaphore_khr clVk2Clsemaphore; - - void* handle1 = NULL; - void* handle2 = NULL; - int fd1, fd2; - std::vector sema_props1{ - (cl_semaphore_properties_khr)CL_SEMAPHORE_TYPE_KHR, - (cl_semaphore_properties_khr)CL_SEMAPHORE_TYPE_BINARY_KHR, - }; - std::vector sema_props2{ - (cl_semaphore_properties_khr)CL_SEMAPHORE_TYPE_KHR, - (cl_semaphore_properties_khr)CL_SEMAPHORE_TYPE_BINARY_KHR, - }; - switch (vkExternalSemaphoreHandleType) + + for (VulkanExternalSemaphoreHandleType semaphoreHandleType : + supportedExternalSemaphores) { + VulkanSemaphore vkVk2Clsemaphore(vkDevice, semaphoreHandleType); + VulkanSemaphore vkCl2Vksemaphore(vkDevice, semaphoreHandleType); + cl_semaphore_khr clCl2Vksemaphore; + cl_semaphore_khr clVk2Clsemaphore; + void* handle1 = NULL; + void* handle2 = NULL; + int fd1, fd2; + std::vector sema_props1{ + (cl_semaphore_properties_khr)CL_SEMAPHORE_TYPE_KHR, + (cl_semaphore_properties_khr)CL_SEMAPHORE_TYPE_BINARY_KHR, + }; + std::vector sema_props2{ + (cl_semaphore_properties_khr)CL_SEMAPHORE_TYPE_KHR, + (cl_semaphore_properties_khr)CL_SEMAPHORE_TYPE_BINARY_KHR, + }; + switch (semaphoreHandleType) + { #ifdef _WIN32 - case VULKAN_EXTERNAL_SEMAPHORE_HANDLE_TYPE_OPAQUE_WIN32_NT: - log_info(" Opaque NT handles are only supported on Windows\n"); - handle1 = vkVk2Clsemaphore.getHandle(vkExternalSemaphoreHandleType); - handle2 = vkCl2Vksemaphore.getHandle(vkExternalSemaphoreHandleType); - errNum = check_external_semaphore_handle_type( - devList[0], CL_SEMAPHORE_HANDLE_OPAQUE_WIN32_KHR); - sema_props1.push_back((cl_semaphore_properties_khr) - CL_SEMAPHORE_HANDLE_OPAQUE_WIN32_KHR); - sema_props1.push_back((cl_semaphore_properties_khr)handle1); - sema_props2.push_back((cl_semaphore_properties_khr) - CL_SEMAPHORE_HANDLE_OPAQUE_WIN32_KHR); - sema_props2.push_back((cl_semaphore_properties_khr)handle2); - break; - case VULKAN_EXTERNAL_SEMAPHORE_HANDLE_TYPE_OPAQUE_WIN32_KMT: - log_info(" Opaque D3DKMT handles are only supported on Windows\n"); - handle1 = vkVk2Clsemaphore.getHandle(vkExternalSemaphoreHandleType); - handle2 = vkCl2Vksemaphore.getHandle(vkExternalSemaphoreHandleType); - errNum = check_external_semaphore_handle_type( - devList[0], CL_SEMAPHORE_HANDLE_OPAQUE_WIN32_KMT_KHR); - sema_props1.push_back((cl_semaphore_properties_khr) - CL_SEMAPHORE_HANDLE_OPAQUE_WIN32_KMT_KHR); - sema_props1.push_back((cl_semaphore_properties_khr)handle1); - sema_props2.push_back((cl_semaphore_properties_khr) - CL_SEMAPHORE_HANDLE_OPAQUE_WIN32_KMT_KHR); - sema_props2.push_back((cl_semaphore_properties_khr)handle2); - break; + case VULKAN_EXTERNAL_SEMAPHORE_HANDLE_TYPE_OPAQUE_WIN32_NT: + log_info(" Opaque NT handles are only supported on Windows\n"); + handle1 = vkVk2Clsemaphore.getHandle(semaphoreHandleType); + handle2 = vkCl2Vksemaphore.getHandle(semaphoreHandleType); + errNum = check_external_semaphore_handle_type( + devList[0], CL_SEMAPHORE_HANDLE_OPAQUE_WIN32_KHR); + sema_props1.push_back((cl_semaphore_properties_khr) + CL_SEMAPHORE_HANDLE_OPAQUE_WIN32_KHR); + sema_props1.push_back((cl_semaphore_properties_khr)handle1); + sema_props2.push_back((cl_semaphore_properties_khr) + CL_SEMAPHORE_HANDLE_OPAQUE_WIN32_KHR); + sema_props2.push_back((cl_semaphore_properties_khr)handle2); + break; + case VULKAN_EXTERNAL_SEMAPHORE_HANDLE_TYPE_OPAQUE_WIN32_KMT: + log_info( + " Opaque D3DKMT handles are only supported on Windows\n"); + handle1 = vkVk2Clsemaphore.getHandle(semaphoreHandleType); + handle2 = vkCl2Vksemaphore.getHandle(semaphoreHandleType); + errNum = check_external_semaphore_handle_type( + devList[0], CL_SEMAPHORE_HANDLE_OPAQUE_WIN32_KMT_KHR); + sema_props1.push_back( + (cl_semaphore_properties_khr) + CL_SEMAPHORE_HANDLE_OPAQUE_WIN32_KMT_KHR); + sema_props1.push_back((cl_semaphore_properties_khr)handle1); + sema_props2.push_back( + (cl_semaphore_properties_khr) + CL_SEMAPHORE_HANDLE_OPAQUE_WIN32_KMT_KHR); + sema_props2.push_back((cl_semaphore_properties_khr)handle2); + break; #else - case VULKAN_EXTERNAL_SEMAPHORE_HANDLE_TYPE_OPAQUE_FD: - log_info(" Opaque file descriptors are not supported on Windows\n"); - fd1 = - (int)vkVk2Clsemaphore.getHandle(vkExternalSemaphoreHandleType); - fd2 = - (int)vkCl2Vksemaphore.getHandle(vkExternalSemaphoreHandleType); - errNum = check_external_semaphore_handle_type( - devList[0], CL_SEMAPHORE_HANDLE_OPAQUE_FD_KHR); - sema_props1.push_back( - (cl_semaphore_properties_khr)CL_SEMAPHORE_HANDLE_OPAQUE_FD_KHR); - sema_props1.push_back((cl_semaphore_properties_khr)fd1); - sema_props2.push_back( - (cl_semaphore_properties_khr)CL_SEMAPHORE_HANDLE_OPAQUE_FD_KHR); - sema_props2.push_back((cl_semaphore_properties_khr)fd2); - break; + case VULKAN_EXTERNAL_SEMAPHORE_HANDLE_TYPE_OPAQUE_FD: + fd1 = (int)vkVk2Clsemaphore.getHandle(semaphoreHandleType); + fd2 = (int)vkCl2Vksemaphore.getHandle(semaphoreHandleType); + errNum = check_external_semaphore_handle_type( + devList[0], CL_SEMAPHORE_HANDLE_OPAQUE_FD_KHR); + sema_props1.push_back((cl_semaphore_properties_khr) + CL_SEMAPHORE_HANDLE_OPAQUE_FD_KHR); + sema_props1.push_back((cl_semaphore_properties_khr)fd1); + sema_props2.push_back((cl_semaphore_properties_khr) + CL_SEMAPHORE_HANDLE_OPAQUE_FD_KHR); + sema_props2.push_back((cl_semaphore_properties_khr)fd2); + break; + case VULKAN_EXTERNAL_SEMAPHORE_HANDLE_TYPE_SYNC_FD: + fd1 = -1; + fd2 = -1; + errNum = check_external_semaphore_handle_type( + devList[0], CL_SEMAPHORE_HANDLE_SYNC_FD_KHR); + sema_props1.push_back((cl_semaphore_properties_khr) + CL_SEMAPHORE_HANDLE_SYNC_FD_KHR); + sema_props1.push_back((cl_semaphore_properties_khr)fd1); + sema_props2.push_back((cl_semaphore_properties_khr) + CL_SEMAPHORE_HANDLE_SYNC_FD_KHR); + sema_props2.push_back((cl_semaphore_properties_khr)fd2); + break; #endif default: log_error("Unsupported external memory handle type\n"); break; - } + } if (CL_SUCCESS != errNum) { throw std::runtime_error( "Unsupported external sempahore handle type\n "); } sema_props1.push_back( - (cl_semaphore_properties_khr)CL_DEVICE_HANDLE_LIST_KHR); + (cl_semaphore_properties_khr)CL_SEMAPHORE_DEVICE_HANDLE_LIST_KHR); sema_props1.push_back((cl_semaphore_properties_khr)devList[0]); sema_props1.push_back( - (cl_semaphore_properties_khr)CL_DEVICE_HANDLE_LIST_END_KHR); + (cl_semaphore_properties_khr)CL_SEMAPHORE_DEVICE_HANDLE_LIST_END_KHR); sema_props2.push_back( - (cl_semaphore_properties_khr)CL_DEVICE_HANDLE_LIST_KHR); + (cl_semaphore_properties_khr)CL_SEMAPHORE_DEVICE_HANDLE_LIST_KHR); sema_props2.push_back((cl_semaphore_properties_khr)devList[0]); sema_props2.push_back( - (cl_semaphore_properties_khr)CL_DEVICE_HANDLE_LIST_END_KHR); + (cl_semaphore_properties_khr)CL_SEMAPHORE_DEVICE_HANDLE_LIST_END_KHR); sema_props1.push_back(0); sema_props2.push_back(0); @@ -532,31 +515,6 @@ int test_consistency_external_semaphore(cl_device_id deviceID, test_error(errNum, "Unable to create semaphore with valid semaphore properties"); - - // Call Signal twice consecutively - errNum = clEnqueueSignalSemaphoresKHRptr(cmd_queue, 1, &clVk2Clsemaphore, - NULL, 0, NULL, NULL); - test_error(errNum, "clEnqueueSignalSemaphoresKHRptr failed"); - - errNum = clEnqueueSignalSemaphoresKHRptr(cmd_queue, 1, &clCl2Vksemaphore, - NULL, 0, NULL, NULL); - test_error(errNum, - "clEnqueueSignalSemaphoresKHRptr failed for two " - "consecutive wait events"); - - - // Call Wait twice consecutively - errNum = clEnqueueWaitSemaphoresKHRptr(cmd_queue, 1, &clVk2Clsemaphore, - NULL, 0, NULL, NULL); - test_error(errNum, "clEnqueueWaitSemaphoresKHRptr failed"); - - errNum = clEnqueueWaitSemaphoresKHRptr(cmd_queue, 1, &clCl2Vksemaphore, - NULL, 0, NULL, NULL); - test_error(errNum, - "clEnqueueWaitSemaphoresKHRptr failed for two " - " consecutive wait events"); - - // Pass invalid object to release call errNum = clReleaseSemaphoreKHRptr(NULL); test_failure_error(errNum, CL_INVALID_VALUE, @@ -569,6 +527,7 @@ int test_consistency_external_semaphore(cl_device_id deviceID, errNum = clReleaseSemaphoreKHRptr(clCl2Vksemaphore); test_error(errNum, "clReleaseSemaphoreKHRptr failed"); + } return TEST_PASS; } diff --git a/test_conformance/vulkan/test_vulkan_interop_buffer.cpp b/test_conformance/vulkan/test_vulkan_interop_buffer.cpp index 559625d718..196a8f33b0 100644 --- a/test_conformance/vulkan/test_vulkan_interop_buffer.cpp +++ b/test_conformance/vulkan/test_vulkan_interop_buffer.cpp @@ -15,15 +15,15 @@ // #include -#include #include #include -#include #include #include +#include #include #include #include "harness/errorHelpers.h" +#include "deviceInfo.h" #define MAX_BUFFERS 5 #define MAX_IMPORTS 5 @@ -80,11 +80,12 @@ __kernel void checkKernel(__global unsigned char *ptr, int size, int expVal, __g } \n\ }"; -int run_test_with_two_queue(cl_context &context, cl_command_queue &cmd_queue1, - cl_command_queue &cmd_queue2, cl_kernel *kernel, - cl_kernel &verify_kernel, VulkanDevice &vkDevice, - uint32_t numBuffers, uint32_t bufferSize, - bool use_fence) +int run_test_with_two_queue( + cl_context &context, cl_command_queue &cmd_queue1, + cl_command_queue &cmd_queue2, cl_kernel *kernel, cl_kernel &verify_kernel, + VulkanDevice &vkDevice, uint32_t numBuffers, uint32_t bufferSize, + bool use_fence, + VulkanExternalSemaphoreHandleType vkExternalSemaphoreHandleType) { int err = CL_SUCCESS; size_t global_work_size[1]; @@ -99,24 +100,15 @@ int run_test_with_two_queue(cl_context &context, cl_command_queue &cmd_queue1, cl_program program = clCreateProgramWithSource( context, 1, &program_source_const, &program_source_length, &err); err = clBuildProgram(program, 0, NULL, NULL, NULL, NULL); - if (err != CL_SUCCESS) - { - print_error(err, "Error: Failed to build program \n"); - return err; - } + test_error(err, "Error: Failed to build program \n"); + // create the kernel kernel_cq = clCreateKernel(program, "clUpdateBuffer", &err); - if (err != CL_SUCCESS) - { - print_error(err, "clCreateKernel failed \n"); - return err; - } + test_error(err, "clCreateKernel failed \n"); const std::vector vkExternalMemoryHandleTypeList = getSupportedVulkanExternalMemoryHandleTypeList(); - VulkanExternalSemaphoreHandleType vkExternalSemaphoreHandleType = - getSupportedVulkanExternalSemaphoreHandleTypeList()[0]; VulkanSemaphore vkVk2CLSemaphore(vkDevice, vkExternalSemaphoreHandleType); VulkanSemaphore vkCl2VkSemaphore(vkDevice, vkExternalSemaphoreHandleType); std::shared_ptr fence = nullptr; @@ -166,6 +158,7 @@ int run_test_with_two_queue(cl_context &context, cl_command_queue &cmd_queue1, vkParamsDeviceMemory.bindBuffer(vkParamsBuffer); std::vector vkBufferListDeviceMemory; std::vector externalMemory; + for (size_t emhtIdx = 0; emhtIdx < vkExternalMemoryHandleTypeList.size(); emhtIdx++) { @@ -197,7 +190,7 @@ int run_test_with_two_queue(cl_context &context, cl_command_queue &cmd_queue1, vkExternalMemoryHandleType)); externalMemory.push_back(new clExternalMemory( vkBufferListDeviceMemory[bIdx], vkExternalMemoryHandleType, - 0, bufferSize, context, deviceId)); + bufferSize, context, deviceId)); } cl_mem buffers[MAX_BUFFERS]; clFinish(cmd_queue1); @@ -258,7 +251,10 @@ int run_test_with_two_queue(cl_context &context, cl_command_queue &cmd_queue1, vkVk2CLSemaphore); } - clVk2CLExternalSemaphore->wait(cmd_queue1); + err = clVk2CLExternalSemaphore->wait(cmd_queue1); + test_error_and_cleanup( + err, CLEANUP, + "Error: failed to wait on CL external semaphore\n"); } @@ -279,36 +275,27 @@ int run_test_with_two_queue(cl_context &context, cl_command_queue &cmd_queue1, err |= clSetKernelArg(kernel_cq, 2, sizeof(cl_mem), (void *)&(buffers[vkBufferList.size() - 1])); + test_error_and_cleanup( + err, CLEANUP, + "Error: Failed to set arg values for kernel\n"); - if (err != CL_SUCCESS) - { - print_error(err, - "Error: Failed to set arg values for kernel\n"); - goto CLEANUP; - } cl_event first_launch; err = clEnqueueNDRangeKernel(cmd_queue1, update_buffer_kernel, 1, NULL, global_work_size, NULL, 0, NULL, &first_launch); - if (err != CL_SUCCESS) - { - print_error(err, - "Error: Failed to launch update_buffer_kernel," - "error\n"); - goto CLEANUP; - } + test_error_and_cleanup( + err, CLEANUP, + "Error: Failed to launch update_buffer_kernel," + "error\n"); err = clEnqueueNDRangeKernel(cmd_queue2, kernel_cq, 1, NULL, global_work_size, NULL, 1, &first_launch, NULL); - if (err != CL_SUCCESS) - { - print_error(err, - "Error: Failed to launch update_buffer_kernel," - "error\n"); - goto CLEANUP; - } + test_error_and_cleanup( + err, CLEANUP, + "Error: Failed to launch update_buffer_kernel," + "error\n"); if (use_fence) { @@ -319,31 +306,27 @@ int run_test_with_two_queue(cl_context &context, cl_command_queue &cmd_queue1, } else if (!use_fence && iter != (maxIter - 1)) { - clCl2VkExternalSemaphore->signal(cmd_queue2); + err = clCl2VkExternalSemaphore->signal(cmd_queue2); + test_error_and_cleanup(err, CLEANUP, + "Failed to signal CL semaphore\n"); } } error_2 = (uint8_t *)malloc(sizeof(uint8_t)); if (NULL == error_2) { - log_error("Not able to allocate memory\n"); - goto CLEANUP; + test_fail_and_cleanup(err, CLEANUP, + "Not able to allocate memory\n"); } clFinish(cmd_queue2); error_1 = clCreateBuffer(context, CL_MEM_WRITE_ONLY, sizeof(uint8_t), NULL, &err); - if (CL_SUCCESS != err) - { - print_error(err, "Error: clCreateBuffer \n"); - goto CLEANUP; - } + test_error_and_cleanup(err, CLEANUP, "Error: clCreateBuffer \n"); + uint8_t val = 0; err = clEnqueueWriteBuffer(cmd_queue1, error_1, CL_TRUE, 0, sizeof(uint8_t), &val, 0, NULL, NULL); - if (err != CL_SUCCESS) - { - print_error(err, "Error: Failed read output, error\n"); - goto CLEANUP; - } + test_error_and_cleanup(err, CLEANUP, + "Error: Failed read output, error\n"); int calc_max_iter; for (int i = 0; i < vkBufferList.size(); i++) @@ -360,36 +343,28 @@ int run_test_with_two_queue(cl_context &context, cl_command_queue &cmd_queue1, &calc_max_iter); err |= clSetKernelArg(verify_kernel, 3, sizeof(cl_mem), (void *)&error_1); - if (err != CL_SUCCESS) - { - print_error(err, - "Error: Failed to set arg values for " - "verify_kernel \n"); - goto CLEANUP; - } + test_error_and_cleanup(err, CLEANUP, + "Error: Failed to set arg values for " + "verify_kernel \n"); + err = clEnqueueNDRangeKernel(cmd_queue1, verify_kernel, 1, NULL, global_work_size, NULL, 0, NULL, NULL); + test_error_and_cleanup(err, CLEANUP, + "Error: Failed to launch verify_kernel," + "error \n"); - if (err != CL_SUCCESS) - { - print_error(err, - "Error: Failed to launch verify_kernel," - "error \n"); - goto CLEANUP; - } err = clEnqueueReadBuffer(cmd_queue1, error_1, CL_TRUE, 0, sizeof(uint8_t), error_2, 0, NULL, NULL); - if (err != CL_SUCCESS) - { - print_error(err, "Error: Failed read output, error \n "); - goto CLEANUP; - } + test_error_and_cleanup(err, CLEANUP, + "Error: Failed read output, error \n"); + if (*error_2 == 1) { - log_error("&&&& vulkan_opencl_buffer test FAILED\n"); - goto CLEANUP; + test_fail_and_cleanup( + err, CLEANUP, + "&&&& vulkan_opencl_buffer test FAILED\n"); } } for (size_t i = 0; i < vkBufferList.size(); i++) @@ -429,10 +404,12 @@ int run_test_with_two_queue(cl_context &context, cl_command_queue &cmd_queue1, return err; } -int run_test_with_one_queue(cl_context &context, cl_command_queue &cmd_queue1, - cl_kernel *kernel, cl_kernel &verify_kernel, - VulkanDevice &vkDevice, uint32_t numBuffers, - uint32_t bufferSize, bool use_fence) +int run_test_with_one_queue( + cl_context &context, cl_command_queue &cmd_queue1, cl_kernel *kernel, + cl_kernel &verify_kernel, VulkanDevice &vkDevice, uint32_t numBuffers, + uint32_t bufferSize, + VulkanExternalSemaphoreHandleType vkExternalSemaphoreHandleType, + bool use_fence) { log_info("RUNNING TEST WITH ONE QUEUE...... \n\n"); size_t global_work_size[1]; @@ -446,8 +423,6 @@ int run_test_with_one_queue(cl_context &context, cl_command_queue &cmd_queue1, const std::vector vkExternalMemoryHandleTypeList = getSupportedVulkanExternalMemoryHandleTypeList(); - VulkanExternalSemaphoreHandleType vkExternalSemaphoreHandleType = - getSupportedVulkanExternalSemaphoreHandleTypeList()[0]; VulkanSemaphore vkVk2CLSemaphore(vkDevice, vkExternalSemaphoreHandleType); VulkanSemaphore vkCl2VkSemaphore(vkDevice, vkExternalSemaphoreHandleType); std::shared_ptr fence = nullptr; @@ -528,7 +503,7 @@ int run_test_with_one_queue(cl_context &context, cl_command_queue &cmd_queue1, vkExternalMemoryHandleType)); externalMemory.push_back(new clExternalMemory( vkBufferListDeviceMemory[bIdx], vkExternalMemoryHandleType, - 0, bufferSize, context, deviceId)); + bufferSize, context, deviceId)); } cl_mem buffers[4]; clFinish(cmd_queue1); @@ -566,6 +541,10 @@ int run_test_with_one_queue(cl_context &context, cl_command_queue &cmd_queue1, { update_buffer_kernel = kernel[2]; } + else + { + test_fail_and_cleanup(err, CLEANUP, "Buffer list size invalid"); + } // global work size should be less than or equal to // bufferSizeList[i] @@ -602,23 +581,18 @@ int run_test_with_one_queue(cl_context &context, cl_command_queue &cmd_queue1, clSetKernelArg(update_buffer_kernel, i + 1, sizeof(cl_mem), (void *)&(buffers[i])); } + test_error_and_cleanup( + err, CLEANUP, + "Error: Failed to set arg values for kernel\n"); - if (err != CL_SUCCESS) - { - print_error(err, - "Error: Failed to set arg values for kernel\n"); - goto CLEANUP; - } err = clEnqueueNDRangeKernel(cmd_queue1, update_buffer_kernel, 1, NULL, global_work_size, NULL, 0, NULL, NULL); - if (err != CL_SUCCESS) - { - print_error(err, - "Error: Failed to launch update_buffer_kernel," - " error\n"); - goto CLEANUP; - } + test_error_and_cleanup( + err, CLEANUP, + "Error: Failed to launch update_buffer_kernel," + " error\n"); + if (use_fence) { clFlush(cmd_queue1); @@ -626,31 +600,27 @@ int run_test_with_one_queue(cl_context &context, cl_command_queue &cmd_queue1, } else if (!use_fence && (iter != (maxIter - 1))) { - clCl2VkExternalSemaphore->signal(cmd_queue1); + err = clCl2VkExternalSemaphore->signal(cmd_queue1); + test_error_and_cleanup(err, CLEANUP, + "Failed to signal CL semaphore\n"); } } error_2 = (uint8_t *)malloc(sizeof(uint8_t)); if (NULL == error_2) { - log_error("Not able to allocate memory\n"); - goto CLEANUP; + test_fail_and_cleanup(err, CLEANUP, + "Not able to allocate memory\n"); } error_1 = clCreateBuffer(context, CL_MEM_WRITE_ONLY, sizeof(uint8_t), NULL, &err); - if (CL_SUCCESS != err) - { - print_error(err, "Error: clCreateBuffer \n"); - goto CLEANUP; - } + test_error_and_cleanup(err, CLEANUP, "Error: clCreateBuffer \n"); + uint8_t val = 0; err = clEnqueueWriteBuffer(cmd_queue1, error_1, CL_TRUE, 0, sizeof(uint8_t), &val, 0, NULL, NULL); - if (CL_SUCCESS != err) - { - print_error(err, "Error: clEnqueueWriteBuffer \n"); - goto CLEANUP; - } + test_error_and_cleanup(err, CLEANUP, + "Error: clEnqueueWriteBuffer \n"); int calc_max_iter = (maxIter * 2); for (int i = 0; i < vkBufferList.size(); i++) @@ -663,35 +633,27 @@ int run_test_with_one_queue(cl_context &context, cl_command_queue &cmd_queue1, &calc_max_iter); err |= clSetKernelArg(verify_kernel, 3, sizeof(cl_mem), (void *)&error_1); - if (err != CL_SUCCESS) - { - print_error( - err, - "Error: Failed to set arg values for verify_kernel \n"); - goto CLEANUP; - } + test_error_and_cleanup( + err, CLEANUP, + "Error: Failed to set arg values for verify_kernel \n"); + err = clEnqueueNDRangeKernel(cmd_queue1, verify_kernel, 1, NULL, global_work_size, NULL, 0, NULL, NULL); - if (err != CL_SUCCESS) - { - print_error( - err, "Error: Failed to launch verify_kernel, error\n"); - goto CLEANUP; - } + test_error_and_cleanup( + err, CLEANUP, + "Error: Failed to launch verify_kernel, error\n"); err = clEnqueueReadBuffer(cmd_queue1, error_1, CL_TRUE, 0, sizeof(uint8_t), error_2, 0, NULL, NULL); - if (err != CL_SUCCESS) - { - print_error(err, "Error: Failed read output, error \n"); - goto CLEANUP; - } + test_error_and_cleanup(err, CLEANUP, + "Error: Failed read output, error \n"); if (*error_2 == 1) { - log_error("&&&& vulkan_opencl_buffer test FAILED\n"); - goto CLEANUP; + test_fail_and_cleanup( + err, CLEANUP, + "&&&& vulkan_opencl_buffer test FAILED\n"); } } for (size_t i = 0; i < vkBufferList.size(); i++) @@ -733,25 +695,22 @@ int run_test_with_one_queue(cl_context &context, cl_command_queue &cmd_queue1, int run_test_with_multi_import_same_ctx( cl_context &context, cl_command_queue &cmd_queue1, cl_kernel *kernel, cl_kernel &verify_kernel, VulkanDevice &vkDevice, uint32_t numBuffers, - uint32_t bufferSize, uint32_t bufferSizeForOffset, float use_fence) + uint32_t bufferSize, bool use_fence, + VulkanExternalSemaphoreHandleType vkExternalSemaphoreHandleType) { size_t global_work_size[1]; uint8_t *error_2; cl_mem error_1; int numImports = numBuffers; - cl_kernel update_buffer_kernel[MAX_IMPORTS]; + cl_kernel update_buffer_kernel; clExternalSemaphore *clVk2CLExternalSemaphore = NULL; clExternalSemaphore *clCl2VkExternalSemaphore = NULL; int err = CL_SUCCESS; int calc_max_iter; - bool withOffset; - uint32_t pBufferSize; const std::vector vkExternalMemoryHandleTypeList = getSupportedVulkanExternalMemoryHandleTypeList(); - VulkanExternalSemaphoreHandleType vkExternalSemaphoreHandleType = - getSupportedVulkanExternalSemaphoreHandleTypeList()[0]; VulkanSemaphore vkVk2CLSemaphore(vkDevice, vkExternalSemaphoreHandleType); VulkanSemaphore vkCl2VkSemaphore(vkDevice, vkExternalSemaphoreHandleType); std::shared_ptr fence = nullptr; @@ -823,50 +782,25 @@ int run_test_with_multi_import_same_ctx( log_info("Memory type index: %d\n", (uint32_t)memoryType); log_info("Memory type property: %d\n", memoryType.getMemoryTypeProperty()); - for (unsigned int withOffset = 0; - withOffset <= (unsigned int)enableOffset; withOffset++) - { - log_info("Running withOffset case %d\n", (uint32_t)withOffset); - if (withOffset) - { - pBufferSize = bufferSizeForOffset; - } - else - { - pBufferSize = bufferSize; - } + + cl_mem buffers[MAX_BUFFERS][MAX_IMPORTS]; - VulkanBufferList vkBufferList(numBuffers, vkDevice, pBufferSize, + VulkanBufferList vkBufferList(numBuffers, vkDevice, bufferSize, vkExternalMemoryHandleType); - uint32_t interBufferOffset = - (uint32_t)(vkBufferList[0].getSize()); for (size_t bIdx = 0; bIdx < numBuffers; bIdx++) { - if (withOffset == 0) - { - vkBufferListDeviceMemory.push_back( - new VulkanDeviceMemory(vkDevice, vkBufferList[bIdx], - memoryType, - vkExternalMemoryHandleType)); - } - if (withOffset == 1) - { - uint32_t totalSize = - (uint32_t)(vkBufferList.size() * interBufferOffset); - vkBufferListDeviceMemory.push_back( - new VulkanDeviceMemory(vkDevice, totalSize, - memoryType, - vkExternalMemoryHandleType)); - } + vkBufferListDeviceMemory.push_back(new VulkanDeviceMemory( + vkDevice, vkBufferList[bIdx], memoryType, + vkExternalMemoryHandleType)); + std::vector pExternalMemory; for (size_t cl_bIdx = 0; cl_bIdx < numImports; cl_bIdx++) { pExternalMemory.push_back(new clExternalMemory( vkBufferListDeviceMemory[bIdx], - vkExternalMemoryHandleType, - withOffset * bIdx * interBufferOffset, pBufferSize, - context, deviceId)); + vkExternalMemoryHandleType, bufferSize, context, + deviceId)); } externalMemory.push_back(pExternalMemory); } @@ -874,16 +808,15 @@ int run_test_with_multi_import_same_ctx( clFinish(cmd_queue1); Params *params = (Params *)vkParamsDeviceMemory.map(); params->numBuffers = numBuffers; - params->bufferSize = pBufferSize; - params->interBufferOffset = interBufferOffset * withOffset; + params->bufferSize = bufferSize; + params->interBufferOffset = 0; vkParamsDeviceMemory.unmap(); vkDescriptorSet.update(0, vkParamsBuffer); for (size_t bIdx = 0; bIdx < vkBufferList.size(); bIdx++) { size_t buffer_size = vkBufferList[bIdx].getSize(); vkBufferListDeviceMemory[bIdx]->bindBuffer( - vkBufferList[bIdx], - bIdx * interBufferOffset * withOffset); + vkBufferList[bIdx], 0); for (size_t cl_bIdx = 0; cl_bIdx < numImports; cl_bIdx++) { buffers[bIdx][cl_bIdx] = @@ -898,15 +831,13 @@ int run_test_with_multi_import_same_ctx( vkComputePipeline, vkPipelineLayout, vkDescriptorSet); vkCommandBuffer.dispatch(512, 1, 1); vkCommandBuffer.end(); - for (int i = 0; i < numImports; i++) - { - update_buffer_kernel[i] = (numBuffers == 1) - ? kernel[0] - : ((numBuffers == 2) ? kernel[1] : kernel[2]); - } + + update_buffer_kernel = (numBuffers == 1) + ? kernel[0] + : ((numBuffers == 2) ? kernel[1] : kernel[2]); // global work size should be less than or equal to // bufferSizeList[i] - global_work_size[0] = pBufferSize; + global_work_size[0] = bufferSize; for (uint32_t iter = 0; iter < maxIter; iter++) { @@ -935,40 +866,36 @@ int run_test_with_multi_import_same_ctx( } else { - clVk2CLExternalSemaphore->wait(cmd_queue1); + err = clVk2CLExternalSemaphore->wait(cmd_queue1); + test_error_and_cleanup( + err, CLEANUP, + "Error: failed to wait on CL external semaphore\n"); } for (uint8_t launchIter = 0; launchIter < numImports; launchIter++) { - err = clSetKernelArg(update_buffer_kernel[launchIter], - 0, sizeof(uint32_t), - (void *)&pBufferSize); + err = clSetKernelArg(update_buffer_kernel, 0, + sizeof(uint32_t), + (void *)&bufferSize); for (int i = 0; i < numBuffers; i++) { err |= clSetKernelArg( - update_buffer_kernel[launchIter], i + 1, - sizeof(cl_mem), + update_buffer_kernel, i + 1, sizeof(cl_mem), (void *)&(buffers[i][launchIter])); } + test_error_and_cleanup( + err, CLEANUP, + "Error: Failed to set arg values for " + "kernel\n "); - if (err != CL_SUCCESS) - { - print_error(err, - "Error: Failed to set arg values for " - "kernel\n "); - goto CLEANUP; - } err = clEnqueueNDRangeKernel( - cmd_queue1, update_buffer_kernel[launchIter], 1, - NULL, global_work_size, NULL, 0, NULL, NULL); - if (err != CL_SUCCESS) - { - print_error(err, - "Error: Failed to launch " - "update_buffer_kernel, error\n "); - goto CLEANUP; - } + cmd_queue1, update_buffer_kernel, 1, NULL, + global_work_size, NULL, 0, NULL, NULL); + test_error_and_cleanup( + err, CLEANUP, + "Error: Failed to launch " + "update_buffer_kernel, error\n "); } if (use_fence) { @@ -976,74 +903,66 @@ int run_test_with_multi_import_same_ctx( } else if (!use_fence && iter != (maxIter - 1)) { - clCl2VkExternalSemaphore->signal(cmd_queue1); + err = clCl2VkExternalSemaphore->signal(cmd_queue1); + test_error_and_cleanup( + err, CLEANUP, "Failed to signal CL semaphore\n"); } } + error_2 = (uint8_t *)malloc(sizeof(uint8_t)); if (NULL == error_2) { - log_error("Not able to allocate memory\n"); - goto CLEANUP; + test_fail_and_cleanup(err, CLEANUP, + "Not able to allocate memory\n"); } error_1 = clCreateBuffer(context, CL_MEM_WRITE_ONLY, sizeof(uint8_t), NULL, &err); - if (CL_SUCCESS != err) - { - print_error(err, "Error: clCreateBuffer \n"); - goto CLEANUP; - } + test_error_and_cleanup(err, CLEANUP, + "Error: clCreateBuffer \n"); + uint8_t val = 0; err = clEnqueueWriteBuffer(cmd_queue1, error_1, CL_TRUE, 0, sizeof(uint8_t), &val, 0, NULL, NULL); - if (CL_SUCCESS != err) - { - print_error(err, "Error: clEnqueueWriteBuffer \n"); - goto CLEANUP; - } - calc_max_iter = maxIter * (numBuffers + 1); + test_error_and_cleanup(err, CLEANUP, + "Error: clEnqueueWriteBuffer \n"); + + calc_max_iter = maxIter * (numImports + 1); for (int i = 0; i < vkBufferList.size(); i++) { err = clSetKernelArg(verify_kernel, 0, sizeof(cl_mem), (void *)&(buffers[i][0])); err |= clSetKernelArg(verify_kernel, 1, sizeof(int), - &pBufferSize); + &bufferSize); err |= clSetKernelArg(verify_kernel, 2, sizeof(int), &calc_max_iter); err |= clSetKernelArg(verify_kernel, 3, sizeof(cl_mem), (void *)&error_1); - if (err != CL_SUCCESS) - { - print_error(err, - "Error: Failed to set arg values for " - "verify_kernel \n"); - goto CLEANUP; - } + test_error_and_cleanup( + err, CLEANUP, + "Error: Failed to set arg values for " + "verify_kernel \n"); + err = clEnqueueNDRangeKernel(cmd_queue1, verify_kernel, 1, NULL, global_work_size, NULL, 0, NULL, NULL); - if (err != CL_SUCCESS) - { - print_error( - err, - "Error: Failed to launch verify_kernel, error\n"); - goto CLEANUP; - } + test_error_and_cleanup( + err, CLEANUP, + "Error: Failed to launch verify_kernel, error\n"); err = clEnqueueReadBuffer(cmd_queue1, error_1, CL_TRUE, 0, sizeof(uint8_t), error_2, 0, NULL, NULL); - if (err != CL_SUCCESS) - { - print_error(err, "Error: Failed read output, error \n"); - goto CLEANUP; - } + test_error_and_cleanup( + err, CLEANUP, "Error: Failed read output, error \n"); + if (*error_2 == 1) { - log_error("&&&& vulkan_opencl_buffer test FAILED\n"); - goto CLEANUP; + test_fail_and_cleanup( + err, CLEANUP, + " vulkan_opencl_buffer test FAILED\n"); } } for (size_t i = 0; i < vkBufferList.size(); i++) @@ -1066,7 +985,6 @@ int run_test_with_multi_import_same_ctx( + numBuffers); } externalMemory.clear(); - } } } CLEANUP: @@ -1103,8 +1021,8 @@ int run_test_with_multi_import_diff_ctx( cl_context &context, cl_context &context2, cl_command_queue &cmd_queue1, cl_command_queue &cmd_queue2, cl_kernel *kernel1, cl_kernel *kernel2, cl_kernel &verify_kernel, cl_kernel verify_kernel2, VulkanDevice &vkDevice, - uint32_t numBuffers, uint32_t bufferSize, uint32_t bufferSizeForOffset, - float use_fence) + uint32_t numBuffers, uint32_t bufferSize, bool use_fence, + VulkanExternalSemaphoreHandleType vkExternalSemaphoreHandleType) { size_t global_work_size[1]; uint8_t *error_3; @@ -1125,8 +1043,6 @@ int run_test_with_multi_import_diff_ctx( const std::vector vkExternalMemoryHandleTypeList = getSupportedVulkanExternalMemoryHandleTypeList(); - VulkanExternalSemaphoreHandleType vkExternalSemaphoreHandleType = - getSupportedVulkanExternalSemaphoreHandleTypeList()[0]; VulkanSemaphore vkVk2CLSemaphore(vkDevice, vkExternalSemaphoreHandleType); VulkanSemaphore vkCl2VkSemaphore(vkDevice, vkExternalSemaphoreHandleType); std::shared_ptr fence = nullptr; @@ -1203,201 +1119,166 @@ int run_test_with_multi_import_diff_ctx( log_info("Memory type property: %d\n", memoryType.getMemoryTypeProperty()); - for (unsigned int withOffset = 0; - withOffset <= (unsigned int)enableOffset; withOffset++) + cl_mem buffers1[MAX_BUFFERS][MAX_IMPORTS]; + cl_mem buffers2[MAX_BUFFERS][MAX_IMPORTS]; + pBufferSize = bufferSize; + VulkanBufferList vkBufferList(numBuffers, vkDevice, pBufferSize, + vkExternalMemoryHandleType); + uint32_t interBufferOffset = (uint32_t)(vkBufferList[0].getSize()); + + for (size_t bIdx = 0; bIdx < numBuffers; bIdx++) { - log_info("Running withOffset case %d\n", (uint32_t)withOffset); - cl_mem buffers1[MAX_BUFFERS][MAX_IMPORTS]; - cl_mem buffers2[MAX_BUFFERS][MAX_IMPORTS]; - if (withOffset) - { - pBufferSize = bufferSizeForOffset; - } - else + vkBufferListDeviceMemory.push_back(new VulkanDeviceMemory( + vkDevice, vkBufferList[bIdx], memoryType, + vkExternalMemoryHandleType)); + std::vector pExternalMemory1; + std::vector pExternalMemory2; + for (size_t cl_bIdx = 0; cl_bIdx < numImports; cl_bIdx++) { - pBufferSize = bufferSize; + pExternalMemory1.push_back( + new clExternalMemory(vkBufferListDeviceMemory[bIdx], + vkExternalMemoryHandleType, + pBufferSize, context, deviceId)); + pExternalMemory2.push_back( + new clExternalMemory(vkBufferListDeviceMemory[bIdx], + vkExternalMemoryHandleType, + pBufferSize, context2, deviceId)); } - VulkanBufferList vkBufferList(numBuffers, vkDevice, pBufferSize, - vkExternalMemoryHandleType); - uint32_t interBufferOffset = - (uint32_t)(vkBufferList[0].getSize()); + externalMemory1.push_back(pExternalMemory1); + externalMemory2.push_back(pExternalMemory2); + } - for (size_t bIdx = 0; bIdx < numBuffers; bIdx++) + clFinish(cmd_queue1); + Params *params = (Params *)vkParamsDeviceMemory.map(); + params->numBuffers = numBuffers; + params->bufferSize = pBufferSize; + vkParamsDeviceMemory.unmap(); + vkDescriptorSet.update(0, vkParamsBuffer); + for (size_t bIdx = 0; bIdx < vkBufferList.size(); bIdx++) + { + size_t buffer_size = vkBufferList[bIdx].getSize(); + vkBufferListDeviceMemory[bIdx]->bindBuffer(vkBufferList[bIdx], + 0); + for (size_t cl_bIdx = 0; cl_bIdx < numImports; cl_bIdx++) { - if (withOffset == 0) - { - vkBufferListDeviceMemory.push_back( - new VulkanDeviceMemory(vkDevice, pBufferSize, - memoryType, - vkExternalMemoryHandleType)); - } - if (withOffset == 1) - { - uint32_t totalSize = - (uint32_t)(vkBufferList.size() * interBufferOffset); - vkBufferListDeviceMemory.push_back( - new VulkanDeviceMemory(vkDevice, totalSize, - memoryType, - vkExternalMemoryHandleType)); - } - std::vector pExternalMemory1; - std::vector pExternalMemory2; - for (size_t cl_bIdx = 0; cl_bIdx < numImports; cl_bIdx++) - { - pExternalMemory1.push_back(new clExternalMemory( - vkBufferListDeviceMemory[bIdx], - vkExternalMemoryHandleType, - withOffset * bIdx * interBufferOffset, pBufferSize, - context, deviceId)); - pExternalMemory2.push_back(new clExternalMemory( - vkBufferListDeviceMemory[bIdx], - vkExternalMemoryHandleType, - withOffset * bIdx * interBufferOffset, pBufferSize, - context2, deviceId)); - } - externalMemory1.push_back(pExternalMemory1); - externalMemory2.push_back(pExternalMemory2); + buffers1[bIdx][cl_bIdx] = externalMemory1[bIdx][cl_bIdx] + ->getExternalMemoryBuffer(); + buffers2[bIdx][cl_bIdx] = externalMemory2[bIdx][cl_bIdx] + ->getExternalMemoryBuffer(); } + vkDescriptorSet.update((uint32_t)bIdx + 1, vkBufferList[bIdx]); + } - clFinish(cmd_queue1); - Params *params = (Params *)vkParamsDeviceMemory.map(); - params->numBuffers = numBuffers; - params->bufferSize = pBufferSize; - params->interBufferOffset = interBufferOffset * withOffset; - vkParamsDeviceMemory.unmap(); - vkDescriptorSet.update(0, vkParamsBuffer); - for (size_t bIdx = 0; bIdx < vkBufferList.size(); bIdx++) - { - size_t buffer_size = vkBufferList[bIdx].getSize(); - vkBufferListDeviceMemory[bIdx]->bindBuffer( - vkBufferList[bIdx], - bIdx * interBufferOffset * withOffset); - for (size_t cl_bIdx = 0; cl_bIdx < numImports; cl_bIdx++) - { - buffers1[bIdx][cl_bIdx] = - externalMemory1[bIdx][cl_bIdx] - ->getExternalMemoryBuffer(); - buffers2[bIdx][cl_bIdx] = - externalMemory2[bIdx][cl_bIdx] - ->getExternalMemoryBuffer(); - } - vkDescriptorSet.update((uint32_t)bIdx + 1, - vkBufferList[bIdx]); - } + vkCommandBuffer.begin(); + vkCommandBuffer.bindPipeline(vkComputePipeline); + vkCommandBuffer.bindDescriptorSets( + vkComputePipeline, vkPipelineLayout, vkDescriptorSet); + vkCommandBuffer.dispatch(512, 1, 1); + vkCommandBuffer.end(); - vkCommandBuffer.begin(); - vkCommandBuffer.bindPipeline(vkComputePipeline); - vkCommandBuffer.bindDescriptorSets( - vkComputePipeline, vkPipelineLayout, vkDescriptorSet); - vkCommandBuffer.dispatch(512, 1, 1); - vkCommandBuffer.end(); + for (int i = 0; i < numImports; i++) + { + update_buffer_kernel1[i] = (numBuffers == 1) + ? kernel1[0] + : ((numBuffers == 2) ? kernel1[1] : kernel1[2]); + update_buffer_kernel2[i] = (numBuffers == 1) + ? kernel2[0] + : ((numBuffers == 2) ? kernel2[1] : kernel2[2]); + } + + // global work size should be less than or equal + // to bufferSizeList[i] + global_work_size[0] = pBufferSize; - for (int i = 0; i < numImports; i++) + for (uint32_t iter = 0; iter < maxIter; iter++) + { + if (use_fence) { - update_buffer_kernel1[i] = (numBuffers == 1) - ? kernel1[0] - : ((numBuffers == 2) ? kernel1[1] : kernel1[2]); - update_buffer_kernel2[i] = (numBuffers == 1) - ? kernel2[0] - : ((numBuffers == 2) ? kernel2[1] : kernel2[2]); + fence->reset(); + vkQueue.submit(vkCommandBuffer, fence); + fence->wait(); } - - // global work size should be less than or equal - // to bufferSizeList[i] - global_work_size[0] = pBufferSize; - - for (uint32_t iter = 0; iter < maxIter; iter++) + else { - if (use_fence) - { - fence->reset(); - vkQueue.submit(vkCommandBuffer, fence); - fence->wait(); - } - else - { - if (iter == 0) - { - vkQueue.submit(vkCommandBuffer, vkVk2CLSemaphore); - } - else - { - vkQueue.submit(vkCl2VkSemaphore, vkCommandBuffer, - vkVk2CLSemaphore); - } - } - - if (use_fence) + if (iter == 0) { - fence->wait(); + vkQueue.submit(vkCommandBuffer, vkVk2CLSemaphore); } else { - clVk2CLExternalSemaphore->wait(cmd_queue1); + vkQueue.submit(vkCl2VkSemaphore, vkCommandBuffer, + vkVk2CLSemaphore); } + } - for (uint8_t launchIter = 0; launchIter < numImports; - launchIter++) - { - err = clSetKernelArg(update_buffer_kernel1[launchIter], - 0, sizeof(uint32_t), - (void *)&pBufferSize); - for (int i = 0; i < numBuffers; i++) - { - err |= clSetKernelArg( - update_buffer_kernel1[launchIter], i + 1, - sizeof(cl_mem), - (void *)&(buffers1[i][launchIter])); - } + if (use_fence) + { + fence->wait(); + } + else + { + err = clVk2CLExternalSemaphore->wait(cmd_queue1); + test_error_and_cleanup( + err, CLEANUP, + "Error: failed to wait on CL external semaphore\n"); + } - if (err != CL_SUCCESS) - { - print_error(err, - "Error: Failed to set arg values for " - "kernel\n "); - goto CLEANUP; - } - err = clEnqueueNDRangeKernel( - cmd_queue1, update_buffer_kernel1[launchIter], 1, - NULL, global_work_size, NULL, 0, NULL, NULL); - if (err != CL_SUCCESS) - { - print_error(err, - "Error: Failed to launch " - "update_buffer_kernel, error\n"); - goto CLEANUP; - } - } - if (use_fence) - { - clFinish(cmd_queue1); - } - else if (!use_fence && iter != (maxIter - 1)) + for (uint8_t launchIter = 0; launchIter < numImports; + launchIter++) + { + err = + clSetKernelArg(update_buffer_kernel1[launchIter], 0, + sizeof(uint32_t), (void *)&pBufferSize); + for (int i = 0; i < numBuffers; i++) { - clCl2VkExternalSemaphore->signal(cmd_queue1); + err |= clSetKernelArg( + update_buffer_kernel1[launchIter], i + 1, + sizeof(cl_mem), (void *)&(buffers1[i][launchIter])); } + test_error_and_cleanup( + err, CLEANUP, + "Error: Failed to set arg values for " + "kernel\n "); + + err = clEnqueueNDRangeKernel( + cmd_queue1, update_buffer_kernel1[launchIter], 1, NULL, + global_work_size, NULL, 0, NULL, NULL); + test_error_and_cleanup(err, CLEANUP, + "Error: Failed to launch " + "update_buffer_kernel, error\n"); } - clFinish(cmd_queue1); - for (uint32_t iter = 0; iter < maxIter; iter++) + if (use_fence) { - if (use_fence) + clFinish(cmd_queue1); + } + else if (!use_fence && iter != (maxIter - 1)) + { + err = clCl2VkExternalSemaphore->signal(cmd_queue1); + test_error_and_cleanup(err, CLEANUP, + "Failed to signal CL semaphore\n"); + } + } + clFinish(cmd_queue1); + for (uint32_t iter = 0; iter < maxIter; iter++) + { + if (use_fence) + { + fence->reset(); + vkQueue.submit(vkCommandBuffer, fence); + fence->wait(); + } + else + { + if (iter == 0) { - fence->reset(); - vkQueue.submit(vkCommandBuffer, fence); - fence->wait(); + vkQueue.submit(vkCommandBuffer, vkVk2CLSemaphore); } else { - if (iter == 0) - { - vkQueue.submit(vkCommandBuffer, vkVk2CLSemaphore); - } - else - { - vkQueue.submit(vkCl2VkSemaphore, vkCommandBuffer, - vkVk2CLSemaphore); - } + vkQueue.submit(vkCl2VkSemaphore, vkCommandBuffer, + vkVk2CLSemaphore); } + } if (use_fence) { @@ -1405,7 +1286,10 @@ int run_test_with_multi_import_diff_ctx( } else { - clVk2CLExternalSemaphore2->wait(cmd_queue2); + err = clVk2CLExternalSemaphore2->wait(cmd_queue2); + test_error_and_cleanup( + err, CLEANUP, + "Error: failed to wait on CL external semaphore\n"); } for (uint8_t launchIter = 0; launchIter < numImports; @@ -1421,24 +1305,18 @@ int run_test_with_multi_import_diff_ctx( sizeof(cl_mem), (void *)&(buffers2[i][launchIter])); } + test_error_and_cleanup( + err, CLEANUP, + "Error: Failed to set arg values for " + "kernel\n "); - if (err != CL_SUCCESS) - { - print_error(err, - "Error: Failed to set arg values for " - "kernel\n "); - goto CLEANUP; - } err = clEnqueueNDRangeKernel( cmd_queue2, update_buffer_kernel2[launchIter], 1, NULL, global_work_size, NULL, 0, NULL, NULL); - if (err != CL_SUCCESS) - { - print_error(err, - "Error: Failed to launch " - "update_buffer_kernel, error\n "); - goto CLEANUP; - } + test_error_and_cleanup( + err, CLEANUP, + "Error: Failed to launch " + "update_buffer_kernel, error\n "); } if (use_fence) { @@ -1446,49 +1324,41 @@ int run_test_with_multi_import_diff_ctx( } else if (!use_fence && iter != (maxIter - 1)) { - clCl2VkExternalSemaphore2->signal(cmd_queue2); + err = clCl2VkExternalSemaphore2->signal(cmd_queue2); + test_error_and_cleanup( + err, CLEANUP, "Failed to signal CL semaphore\n"); } - } + } clFinish(cmd_queue2); error_3 = (uint8_t *)malloc(sizeof(uint8_t)); if (NULL == error_3) { - log_error("Not able to allocate memory\n"); - goto CLEANUP; + test_fail_and_cleanup(err, CLEANUP, + "Not able to allocate memory\n"); } error_1 = clCreateBuffer(context, CL_MEM_WRITE_ONLY, sizeof(uint8_t), NULL, &err); - if (CL_SUCCESS != err) - { - print_error(err, "Error: clCreateBuffer \n"); - goto CLEANUP; - } + test_error_and_cleanup(err, CLEANUP, + "Error: clCreateBuffer \n"); + error_2 = clCreateBuffer(context2, CL_MEM_WRITE_ONLY, sizeof(uint8_t), NULL, &err); - if (CL_SUCCESS != err) - { - print_error(err, "Error: clCreateBuffer \n"); - goto CLEANUP; - } + test_error_and_cleanup(err, CLEANUP, + "Error: clCreateBuffer \n"); + uint8_t val = 0; err = clEnqueueWriteBuffer(cmd_queue1, error_1, CL_TRUE, 0, sizeof(uint8_t), &val, 0, NULL, NULL); - if (err != CL_SUCCESS) - { - print_error(err, "Error: Failed read output, error \n"); - goto CLEANUP; - } + test_error_and_cleanup(err, CLEANUP, + "Error: Failed read output, error \n"); err = clEnqueueWriteBuffer(cmd_queue2, error_2, CL_TRUE, 0, sizeof(uint8_t), &val, 0, NULL, NULL); - if (err != CL_SUCCESS) - { - print_error(err, "Error: Failed read output, error \n"); - goto CLEANUP; - } + test_error_and_cleanup(err, CLEANUP, + "Error: Failed read output, error \n"); calc_max_iter = maxIter * 2 * (numBuffers + 1); for (int i = 0; i < numBuffers; i++) @@ -1501,36 +1371,30 @@ int run_test_with_multi_import_diff_ctx( &calc_max_iter); err |= clSetKernelArg(verify_kernel, 3, sizeof(cl_mem), (void *)&error_1); - if (err != CL_SUCCESS) - { - print_error(err, - "Error: Failed to set arg values for " - "verify_kernel \n"); - goto CLEANUP; - } + test_error_and_cleanup( + err, CLEANUP, + "Error: Failed to set arg values for " + "verify_kernel \n"); + err = clEnqueueNDRangeKernel(cmd_queue1, verify_kernel, 1, NULL, global_work_size, NULL, 0, NULL, NULL); - if (err != CL_SUCCESS) - { - print_error(err, - "Error: Failed to launch verify_kernel," - "error\n"); - goto CLEANUP; - } + test_error_and_cleanup( + err, CLEANUP, + "Error: Failed to launch verify_kernel," + "error\n"); err = clEnqueueReadBuffer(cmd_queue1, error_1, CL_TRUE, 0, sizeof(uint8_t), error_3, 0, NULL, NULL); - if (err != CL_SUCCESS) - { - print_error(err, "Error: Failed read output, error\n"); - goto CLEANUP; - } + test_error_and_cleanup( + err, CLEANUP, "Error: Failed read output, error\n"); + if (*error_3 == 1) { - log_error("&&&& vulkan_opencl_buffer test FAILED\n"); - goto CLEANUP; + test_fail_and_cleanup( + err, CLEANUP, + "&&&& vulkan_opencl_buffer test FAILED\n"); } } *error_3 = 0; @@ -1544,36 +1408,30 @@ int run_test_with_multi_import_diff_ctx( &calc_max_iter); err |= clSetKernelArg(verify_kernel2, 3, sizeof(cl_mem), (void *)&error_2); - if (err != CL_SUCCESS) - { - print_error(err, - "Error: Failed to set arg values for " - "verify_kernel \n"); - goto CLEANUP; - } + test_error_and_cleanup( + err, CLEANUP, + "Error: Failed to set arg values for " + "verify_kernel \n"); + err = clEnqueueNDRangeKernel(cmd_queue2, verify_kernel2, 1, NULL, global_work_size, NULL, 0, NULL, NULL); - if (err != CL_SUCCESS) - { - print_error(err, - "Error: Failed to launch verify_kernel," - "error\n"); - goto CLEANUP; - } + test_error_and_cleanup( + err, CLEANUP, + "Error: Failed to launch verify_kernel," + "error\n"); err = clEnqueueReadBuffer(cmd_queue2, error_2, CL_TRUE, 0, sizeof(uint8_t), error_3, 0, NULL, NULL); - if (err != CL_SUCCESS) - { - print_error(err, "Error: Failed read output, error\n"); - goto CLEANUP; - } + test_error_and_cleanup( + err, CLEANUP, "Error: Failed read output, error\n"); + if (*error_3 == 1) { - log_error("&&&& vulkan_opencl_buffer test FAILED\n"); - goto CLEANUP; + test_fail_and_cleanup( + err, CLEANUP, + "&&&& vulkan_opencl_buffer test FAILED\n"); } } for (size_t i = 0; i < vkBufferList.size(); i++) @@ -1601,7 +1459,6 @@ int run_test_with_multi_import_diff_ctx( } externalMemory1.clear(); externalMemory2.clear(); - } } } CLEANUP: @@ -1649,7 +1506,7 @@ int run_test_with_multi_import_diff_ctx( int test_buffer_common(cl_device_id device_, cl_context context_, cl_command_queue queue_, int numElements_, - float use_fence) + bool use_fence) { int current_device = 0; @@ -1688,69 +1545,53 @@ int test_buffer_common(cl_device_id device_, cl_context context_, uint32_t bufferSizeListforOffset[] = { 256, 512, 1024 }; cl_context_properties contextProperties[] = { CL_CONTEXT_PLATFORM, 0, 0 }; + std::vector supportedSemaphoreTypes; + errNum = clGetPlatformIDs(1, &platform, NULL); - if (errNum != CL_SUCCESS) - { - print_error(errNum, "Error: Failed to get platform\n"); - goto CLEANUP; - } + test_error_and_cleanup(errNum, CLEANUP, "Error: Failed to get platform\n"); errNum = clGetDeviceIDs(platform, CL_DEVICE_TYPE_GPU, 0, NULL, &num_devices); - if (CL_SUCCESS != errNum) - { - print_error(errNum, "clGetDeviceIDs failed in returning of devices\n"); - goto CLEANUP; - } + test_error_and_cleanup(errNum, CLEANUP, + "clGetDeviceIDs failed in returning of devices\n"); + devices = (cl_device_id *)malloc(num_devices * sizeof(cl_device_id)); if (NULL == devices) { - errNum = CL_OUT_OF_HOST_MEMORY; - print_error(errNum, "Unable to allocate memory for devices\n"); - goto CLEANUP; + test_fail_and_cleanup(errNum, CLEANUP, + "Unable to allocate memory for devices\n"); } errNum = clGetDeviceIDs(platform, CL_DEVICE_TYPE_GPU, num_devices, devices, NULL); - if (CL_SUCCESS != errNum) - { - print_error(errNum, "Failed to get deviceID.\n"); - goto CLEANUP; - } + test_error_and_cleanup(errNum, CLEANUP, "Failed to get deviceID.\n"); + contextProperties[1] = (cl_context_properties)platform; log_info("Assigned contextproperties for platform\n"); for (device_no = 0; device_no < num_devices; device_no++) { - errNum = clGetDeviceInfo(devices[device_no], CL_DEVICE_EXTENSIONS, 0, - NULL, &extensionSize); - if (CL_SUCCESS != errNum) - { - print_error(errNum, - "Error in clGetDeviceInfo for getting device_extension " - "size....\n"); - goto CLEANUP; - } - extensions = (char *)malloc(extensionSize); - if (NULL == extensions) + errNum = clGetDeviceInfo(devices[device_no], CL_DEVICE_UUID_KHR, + CL_UUID_SIZE_KHR, uuid, NULL); + test_error_and_cleanup(errNum, CLEANUP, "clGetDeviceInfo failed\n"); + + if (!use_fence) { - print_error(errNum, "Unable to allocate memory for extensions\n"); - errNum = CL_OUT_OF_HOST_MEMORY; - goto CLEANUP; + supportedSemaphoreTypes = + getSupportedInteropExternalSemaphoreHandleTypes( + devices[device_no], vkDevice); } - errNum = clGetDeviceInfo(devices[device_no], CL_DEVICE_EXTENSIONS, - extensionSize, extensions, NULL); - if (CL_SUCCESS != errNum) + else { - print_error(errNum, - "Error in clGetDeviceInfo for device_extension\n"); - goto CLEANUP; + supportedSemaphoreTypes.push_back( + VULKAN_EXTERNAL_SEMAPHORE_HANDLE_TYPE_NONE); } - errNum = clGetDeviceInfo(devices[device_no], CL_DEVICE_UUID_KHR, - CL_UUID_SIZE_KHR, uuid, &extensionSize); - if (CL_SUCCESS != errNum) + + + // If device does not support any semaphores, try the next one + if (!use_fence && supportedSemaphoreTypes.empty()) { - print_error(errNum, "clGetDeviceInfo failed\n"); - goto CLEANUP; + continue; } + errNum = memcmp(uuid, vkDevice.getPhysicalDevice().getUUID(), VK_UUID_SIZE); if (errNum == 0) @@ -1758,38 +1599,35 @@ int test_buffer_common(cl_device_id device_, cl_context context_, break; } } + + if (!use_fence && supportedSemaphoreTypes.empty()) + { + test_fail_and_cleanup( + errNum, CLEANUP, + "No devices found that support OpenCL semaphores\n"); + } + if (device_no >= num_devices) { - errNum = EXIT_FAILURE; - print_error(errNum, - "OpenCL error: " - "No Vulkan-OpenCL Interop capable GPU found.\n"); - goto CLEANUP; + test_fail_and_cleanup(errNum, CLEANUP, + "OpenCL error: " + "No Vulkan-OpenCL Interop capable GPU found.\n"); } deviceId = devices[device_no]; context = clCreateContextFromType(contextProperties, CL_DEVICE_TYPE_GPU, NULL, NULL, &errNum); - if (CL_SUCCESS != errNum) - { - print_error(errNum, "error creating context\n"); - goto CLEANUP; - } + test_error_and_cleanup(errNum, CLEANUP, "error creating context\n"); + log_info("Successfully created context !!!\n"); cmd_queue1 = clCreateCommandQueue(context, devices[device_no], 0, &errNum); - if (CL_SUCCESS != errNum) - { - errNum = CL_INVALID_COMMAND_QUEUE; - print_error(errNum, "Error: Failed to create command queue!\n"); - goto CLEANUP; - } + test_error_and_cleanup(errNum, CLEANUP, + "Error: Failed to create command queue!\n"); + cmd_queue2 = clCreateCommandQueue(context, devices[device_no], 0, &errNum); - if (CL_SUCCESS != errNum) - { - errNum = CL_INVALID_COMMAND_QUEUE; - print_error(errNum, "Error: Failed to create command queue!\n"); - goto CLEANUP; - } + test_error_and_cleanup(errNum, CLEANUP, + "Error: Failed to create command queue!\n"); + log_info("clCreateCommandQueue successful\n"); for (int i = 0; i < 3; i++) { @@ -1798,18 +1636,12 @@ int test_buffer_common(cl_device_id device_, cl_context context_, clCreateProgramWithSource(context, 1, &program_source_const[i], &program_source_length, &errNum); errNum = clBuildProgram(program[i], 0, NULL, NULL, NULL, NULL); - if (errNum != CL_SUCCESS) - { - print_error(errNum, "Error: Failed to build program \n"); - return errNum; - } + test_error_and_cleanup(errNum, CLEANUP, + "Error: Failed to build program \n"); + // create the kernel kernel[i] = clCreateKernel(program[i], "clUpdateBuffer", &errNum); - if (errNum != CL_SUCCESS) - { - print_error(errNum, "clCreateKernel failed \n"); - return errNum; - } + test_error_and_cleanup(errNum, CLEANUP, "clCreateKernel failed \n"); } program_source_const_verify = kernel_text_verify; @@ -1818,35 +1650,23 @@ int test_buffer_common(cl_device_id device_, cl_context context_, clCreateProgramWithSource(context, 1, &program_source_const_verify, &program_source_length, &errNum); errNum = clBuildProgram(program_verify, 0, NULL, NULL, NULL, NULL); - if (errNum != CL_SUCCESS) - { - log_error("Error: Failed to build program2\n"); - return errNum; - } + test_error_and_cleanup(errNum, CLEANUP, + "Error: Failed to build program2\n"); + verify_kernel = clCreateKernel(program_verify, "checkKernel", &errNum); - if (errNum != CL_SUCCESS) - { - print_error(errNum, "clCreateKernel failed \n"); - return errNum; - } + test_error_and_cleanup(errNum, CLEANUP, "clCreateKernel failed \n"); if (multiCtx) // different context guard { context2 = clCreateContextFromType( contextProperties, CL_DEVICE_TYPE_GPU, NULL, NULL, &errNum); - if (CL_SUCCESS != errNum) - { - print_error(errNum, "error creating context\n"); - goto CLEANUP; - } + test_error_and_cleanup(errNum, CLEANUP, "error creating context\n"); + cmd_queue3 = clCreateCommandQueue(context2, devices[device_no], 0, &errNum); - if (CL_SUCCESS != errNum) - { - errNum = CL_INVALID_COMMAND_QUEUE; - print_error(errNum, "Error: Failed to create command queue!\n"); - goto CLEANUP; - } + test_error_and_cleanup(errNum, CLEANUP, + "Error: Failed to create command queue!\n"); + for (int i = 0; i < 3; i++) { program_source_length = strlen(program_source_const[i]); @@ -1854,79 +1674,69 @@ int test_buffer_common(cl_device_id device_, cl_context context_, clCreateProgramWithSource(context2, 1, &program_source_const[i], &program_source_length, &errNum); errNum = clBuildProgram(program[i], 0, NULL, NULL, NULL, NULL); - if (errNum != CL_SUCCESS) - { - print_error(errNum, "Error: Failed to build program \n"); - return errNum; - } + test_error_and_cleanup(errNum, CLEANUP, + "Error: Failed to build program \n"); + // create the kernel kernel2[i] = clCreateKernel(program[i], "clUpdateBuffer", &errNum); - if (errNum != CL_SUCCESS) - { - print_error(errNum, "clCreateKernel failed \n"); - return errNum; - } + test_error_and_cleanup(errNum, CLEANUP, "clCreateKernel failed \n"); } program_source_length = strlen(program_source_const_verify); program_verify = clCreateProgramWithSource(context2, 1, &program_source_const_verify, &program_source_length, &errNum); errNum = clBuildProgram(program_verify, 0, NULL, NULL, NULL, NULL); - if (errNum != CL_SUCCESS) - { - log_error("Error: Failed to build program2\n"); - return errNum; - } + test_error_and_cleanup(errNum, CLEANUP, + "Error: Failed to build program2\n"); + verify_kernel2 = clCreateKernel(program_verify, "checkKernel", &errNum); - if (errNum != CL_SUCCESS) - { - print_error(errNum, "clCreateKernel failed \n"); - return errNum; - } + test_error_and_cleanup(errNum, CLEANUP, "clCreateKernel failed \n"); } - for (size_t numBuffersIdx = 0; numBuffersIdx < ARRAY_SIZE(numBuffersList); - numBuffersIdx++) + // TODO: Add support for empty list if use_fence enabled + for (VulkanExternalSemaphoreHandleType semaphoreType : + supportedSemaphoreTypes) { - uint32_t numBuffers = numBuffersList[numBuffersIdx]; - log_info("Number of buffers: %d\n", numBuffers); - for (size_t sizeIdx = 0; sizeIdx < ARRAY_SIZE(bufferSizeList); - sizeIdx++) + for (size_t numBuffersIdx = 0; + numBuffersIdx < ARRAY_SIZE(numBuffersList); numBuffersIdx++) { - uint32_t bufferSize = bufferSizeList[sizeIdx]; - uint32_t bufferSizeForOffset = bufferSizeListforOffset[sizeIdx]; - log_info("&&&& RUNNING vulkan_opencl_buffer test for Buffer size: " - "%d\n", - bufferSize); - if (multiImport && !multiCtx) - { - errNum = run_test_with_multi_import_same_ctx( - context, cmd_queue1, kernel, verify_kernel, vkDevice, - numBuffers, bufferSize, bufferSizeForOffset, use_fence); - } - else if (multiImport && multiCtx) - { - errNum = run_test_with_multi_import_diff_ctx( - context, context2, cmd_queue1, cmd_queue3, kernel, kernel2, - verify_kernel, verify_kernel2, vkDevice, numBuffers, - bufferSize, bufferSizeForOffset, use_fence); - } - else if (numCQ == 2) - { - errNum = run_test_with_two_queue( - context, cmd_queue1, cmd_queue2, kernel, verify_kernel, - vkDevice, numBuffers + 1, bufferSize, use_fence); - } - else + uint32_t numBuffers = numBuffersList[numBuffersIdx]; + log_info("Number of buffers: %d\n", numBuffers); + for (size_t sizeIdx = 0; sizeIdx < ARRAY_SIZE(bufferSizeList); + sizeIdx++) { - errNum = run_test_with_one_queue( - context, cmd_queue1, kernel, verify_kernel, vkDevice, - numBuffers, bufferSize, use_fence); - } - if (errNum != CL_SUCCESS) - { - print_error(errNum, "func_name failed \n"); - goto CLEANUP; + uint32_t bufferSize = bufferSizeList[sizeIdx]; + log_info( + "&&&& RUNNING vulkan_opencl_buffer test for Buffer size: " + "%d\n", + bufferSize); + if (multiImport && !multiCtx) + { + errNum = run_test_with_multi_import_same_ctx( + context, cmd_queue1, kernel, verify_kernel, vkDevice, + numBuffers, bufferSize, use_fence, semaphoreType); + } + else if (multiImport && multiCtx) + { + errNum = run_test_with_multi_import_diff_ctx( + context, context2, cmd_queue1, cmd_queue3, kernel, + kernel2, verify_kernel, verify_kernel2, vkDevice, + numBuffers, bufferSize, use_fence, semaphoreType); + } + else if (numCQ == 2) + { + errNum = run_test_with_two_queue( + context, cmd_queue1, cmd_queue2, kernel, verify_kernel, + vkDevice, numBuffers + 1, bufferSize, use_fence, + semaphoreType); + } + else + { + errNum = run_test_with_one_queue( + context, cmd_queue1, kernel, verify_kernel, vkDevice, + numBuffers, bufferSize, semaphoreType, use_fence); + } + test_error_and_cleanup(errNum, CLEANUP, "func_name failed \n"); } } } diff --git a/test_conformance/vulkan/test_vulkan_interop_image.cpp b/test_conformance/vulkan/test_vulkan_interop_image.cpp index 5f1f6e4b41..872044df9d 100644 --- a/test_conformance/vulkan/test_vulkan_interop_image.cpp +++ b/test_conformance/vulkan/test_vulkan_interop_image.cpp @@ -18,6 +18,7 @@ #include #include "harness/errorHelpers.h" #include +#include "deviceInfo.h" #define MAX_2D_IMAGES 5 #define MAX_2D_IMAGE_WIDTH 1024 @@ -189,11 +190,11 @@ const cl_kernel getKernelType(VulkanFormat format, cl_kernel kernel_float, return kernel; } -int run_test_with_two_queue(cl_context &context, cl_command_queue &cmd_queue1, - cl_command_queue &cmd_queue2, - cl_kernel *kernel_unsigned, - cl_kernel *kernel_signed, cl_kernel *kernel_float, - VulkanDevice &vkDevice) +int run_test_with_two_queue( + cl_context &context, cl_command_queue &cmd_queue1, + cl_command_queue &cmd_queue2, cl_kernel *kernel_unsigned, + cl_kernel *kernel_signed, cl_kernel *kernel_float, VulkanDevice &vkDevice, + VulkanExternalSemaphoreHandleType vkExternalSemaphoreHandleType) { cl_int err = CL_SUCCESS; size_t origin[3] = { 0, 0, 0 }; @@ -245,8 +246,6 @@ int run_test_with_two_queue(cl_context &context, cl_command_queue &cmd_queue1, VulkanCommandBuffer vkShaderCommandBuffer(vkDevice, vkCommandPool); VulkanQueue &vkQueue = vkDevice.getQueue(); - VulkanExternalSemaphoreHandleType vkExternalSemaphoreHandleType = - getSupportedVulkanExternalSemaphoreHandleTypeList()[0]; VulkanSemaphore vkVk2CLSemaphore(vkDevice, vkExternalSemaphoreHandleType); VulkanSemaphore vkCl2VkSemaphore(vkDevice, vkExternalSemaphoreHandleType); clExternalSemaphore *clVk2CLExternalSemaphore = NULL; @@ -462,7 +461,11 @@ int run_test_with_two_queue(cl_context &context, cl_command_queue &cmd_queue1, ->getExternalMemoryImage(); } - clCl2VkExternalSemaphore->signal(cmd_queue1); + err = clCl2VkExternalSemaphore->signal(cmd_queue1); + test_error_and_cleanup( + err, CLEANUP, + "Failed to signal CL semaphore\n"); + if (!useSingleImageKernel) { vkDescriptorSet.updateArray(1, @@ -499,6 +502,7 @@ int run_test_with_two_queue(cl_context &context, cl_command_queue &cmd_queue1, 1); vkShaderCommandBuffer.end(); } + for (uint32_t iter = 0; iter < innerIterations; iter++) { @@ -552,7 +556,17 @@ int run_test_with_two_queue(cl_context &context, cl_command_queue &cmd_queue1, vkQueue.submit(vkCl2VkSemaphore, vkShaderCommandBuffer, vkVk2CLSemaphore); - clVk2CLExternalSemaphore->wait(cmd_queue1); + + err = + clVk2CLExternalSemaphore->wait(cmd_queue1); + if (err != CL_SUCCESS) + { + print_error(err, + "Error: failed to wait on CL " + "external semaphore\n"); + goto CLEANUP; + } + switch (num2DImages) { case 2: @@ -626,14 +640,10 @@ int run_test_with_two_queue(cl_context &context, cl_command_queue &cmd_queue1, err |= clSetKernelArg(updateKernelCQ1, ++j, sizeof(unsigned int), &numMipLevels); + test_error_and_cleanup( + err, CLEANUP, + "Error: Failed to set arg values \n"); - if (err != CL_SUCCESS) - { - print_error( - err, - "Error: Failed to set arg values \n"); - goto CLEANUP; - } // clVk2CLExternalSemaphore->wait(cmd_queue1); size_t global_work_size[3] = { width, height, 1 }; @@ -642,21 +652,24 @@ int run_test_with_two_queue(cl_context &context, cl_command_queue &cmd_queue1, cmd_queue1, updateKernelCQ1, 2, NULL, global_work_size, NULL, 0, NULL, &first_launch); - if (err != CL_SUCCESS) - { - goto CLEANUP; - } + test_error_and_cleanup( + err, CLEANUP, + "Failed to enqueue updateKernelCQ1\n"); + err = clEnqueueNDRangeKernel( cmd_queue2, updateKernelCQ2, 2, NULL, global_work_size, NULL, 1, &first_launch, NULL); - if (err != CL_SUCCESS) - { - goto CLEANUP; - } + test_error_and_cleanup( + err, CLEANUP, + "Failed to enqueue updateKernelCQ2\n"); clFinish(cmd_queue2); - clCl2VkExternalSemaphore->signal(cmd_queue2); + err = clCl2VkExternalSemaphore->signal( + cmd_queue2); + test_error_and_cleanup( + err, CLEANUP, + "Failed to signal CL semaphore\n"); } unsigned int flags = 0; @@ -668,14 +681,11 @@ int run_test_with_two_queue(cl_context &context, cl_command_queue &cmd_queue1, err = clEnqueueReadImage( cmd_queue1, external_mem_image2[i], CL_TRUE, origin, region, 0, 0, dstBufferPtr, 0, NULL, - &eventReadImage); - - if (err != CL_SUCCESS) - { - print_error(err, - "clEnqueueReadImage failed with" - "error\n"); - } + NULL); + test_error_and_cleanup( + err, CLEANUP, + "clEnqueueReadImage failed with" + "error\n"); if (memcmp(srcBufferPtr, dstBufferPtr, srcBufSize)) @@ -727,10 +737,8 @@ int run_test_with_two_queue(cl_context &context, cl_command_queue &cmd_queue1, externalMemory2.erase(externalMemory2.begin(), externalMemory2.begin() + num2DImages); - if (CL_SUCCESS != err) - { - goto CLEANUP; - } + test_error_and_cleanup(err, CLEANUP, + "Test error detected\n"); } } } @@ -748,10 +756,11 @@ int run_test_with_two_queue(cl_context &context, cl_command_queue &cmd_queue1, return err; } -int run_test_with_one_queue(cl_context &context, cl_command_queue &cmd_queue1, - cl_kernel *kernel_unsigned, - cl_kernel *kernel_signed, cl_kernel *kernel_float, - VulkanDevice &vkDevice) +int run_test_with_one_queue( + cl_context &context, cl_command_queue &cmd_queue1, + cl_kernel *kernel_unsigned, cl_kernel *kernel_signed, + cl_kernel *kernel_float, VulkanDevice &vkDevice, + VulkanExternalSemaphoreHandleType vkExternalSemaphoreHandleType) { cl_int err = CL_SUCCESS; size_t origin[3] = { 0, 0, 0 }; @@ -802,8 +811,6 @@ int run_test_with_one_queue(cl_context &context, cl_command_queue &cmd_queue1, VulkanCommandBuffer vkShaderCommandBuffer(vkDevice, vkCommandPool); VulkanQueue &vkQueue = vkDevice.getQueue(); - VulkanExternalSemaphoreHandleType vkExternalSemaphoreHandleType = - getSupportedVulkanExternalSemaphoreHandleTypeList()[0]; VulkanSemaphore vkVk2CLSemaphore(vkDevice, vkExternalSemaphoreHandleType); VulkanSemaphore vkCl2VkSemaphore(vkDevice, vkExternalSemaphoreHandleType); clExternalSemaphore *clVk2CLExternalSemaphore = NULL; @@ -925,8 +932,8 @@ int run_test_with_one_queue(cl_context &context, cl_command_queue &cmd_queue1, vkClExternalMemoryHandleTilingAssumption( deviceId, vkExternalMemoryHandleTypeList[emhtIdx], &err); - ASSERT_SUCCESS(err, - "Failed to query OpenCL tiling mode"); + test_error_and_cleanup( + err, CLEANUP, "Failed to query OpenCL tiling mode"); VulkanImage2D vkDummyImage2D( vkDevice, vkFormatList[0], widthList[0], @@ -1024,7 +1031,11 @@ int run_test_with_one_queue(cl_context &context, cl_command_queue &cmd_queue1, ->getExternalMemoryImage(); } - clCl2VkExternalSemaphore->signal(cmd_queue1); + err = clCl2VkExternalSemaphore->signal(cmd_queue1); + test_error_and_cleanup( + err, CLEANUP, + "Failed to signal CL semaphore\n"); + if (!useSingleImageKernel) { vkDescriptorSet.updateArray(1, @@ -1061,6 +1072,7 @@ int run_test_with_one_queue(cl_context &context, cl_command_queue &cmd_queue1, 1); vkShaderCommandBuffer.end(); } + for (uint32_t iter = 0; iter < innerIterations; iter++) { @@ -1114,7 +1126,14 @@ int run_test_with_one_queue(cl_context &context, cl_command_queue &cmd_queue1, vkQueue.submit(vkCl2VkSemaphore, vkShaderCommandBuffer, vkVk2CLSemaphore); - clVk2CLExternalSemaphore->wait(cmd_queue1); + + err = + clVk2CLExternalSemaphore->wait(cmd_queue1); + test_error_and_cleanup( + err, CLEANUP, + "Error: failed to wait on CL external " + "semaphore\n"); + switch (num2DImages) { case 1: @@ -1158,25 +1177,25 @@ int run_test_with_one_queue(cl_context &context, cl_command_queue &cmd_queue1, err |= clSetKernelArg(updateKernelCQ1, ++j, sizeof(unsigned int), &numMipLevels); - - if (err != CL_SUCCESS) - { - print_error(err, - "Error: Failed to set arg " - "values for kernel-1\n"); - goto CLEANUP; - } + test_error_and_cleanup( + err, CLEANUP, + "Error: Failed to set arg " + "values for kernel-1\n"); size_t global_work_size[3] = { width, height, 1 }; err = clEnqueueNDRangeKernel( cmd_queue1, updateKernelCQ1, 2, NULL, global_work_size, NULL, 0, NULL, NULL); - if (err != CL_SUCCESS) - { - goto CLEANUP; - } - clCl2VkExternalSemaphore->signal(cmd_queue1); + test_error_and_cleanup( + err, CLEANUP, + "Failed to enqueue updateKernelCQ1\n"); + + err = clCl2VkExternalSemaphore->signal( + cmd_queue1); + test_error_and_cleanup( + err, CLEANUP, + "Failed to signal CL semaphore\n"); } unsigned int flags = 0; @@ -1187,14 +1206,11 @@ int run_test_with_one_queue(cl_context &context, cl_command_queue &cmd_queue1, err = clEnqueueReadImage( cmd_queue1, external_mem_image2[i], CL_TRUE, origin, region, 0, 0, dstBufferPtr, 0, NULL, - &eventReadImage); - - if (err != CL_SUCCESS) - { - print_error(err, - "clEnqueueReadImage failed with" - "error\n"); - } + NULL); + test_error_and_cleanup( + err, CLEANUP, + "clEnqueueReadImage failed with" + "error\n"); if (memcmp(srcBufferPtr, dstBufferPtr, srcBufSize)) @@ -1246,10 +1262,8 @@ int run_test_with_one_queue(cl_context &context, cl_command_queue &cmd_queue1, externalMemory2.erase(externalMemory2.begin(), externalMemory2.begin() + num2DImages); - if (CL_SUCCESS != err) - { - goto CLEANUP; - } + test_error_and_cleanup(err, CLEANUP, + "Test detected error\n"); } } } @@ -1293,44 +1307,35 @@ int test_image_common(cl_device_id device_, cl_context context_, char source_2[4096]; char source_3[4096]; size_t program_source_length; - cl_program program[num_kernel_types]; - cl_kernel kernel_float[num_kernels] = { NULL, NULL, NULL, NULL }; - cl_kernel kernel_signed[num_kernels] = { NULL, NULL, NULL, NULL }; - cl_kernel kernel_unsigned[num_kernels] = { NULL, NULL, NULL, NULL }; + cl_program program[num_kernel_types] = { NULL }; + cl_kernel kernel_float[num_kernels] = { NULL }; + cl_kernel kernel_signed[num_kernels] = { NULL }; + cl_kernel kernel_unsigned[num_kernels] = { NULL }; cl_mem external_mem_image1; cl_mem external_mem_image2; + std::vector supportedSemaphoreTypes; VulkanDevice vkDevice; cl_context_properties contextProperties[] = { CL_CONTEXT_PLATFORM, 0, 0 }; // get the platform ID err = clGetPlatformIDs(1, &platform, NULL); - if (err != CL_SUCCESS) - { - print_error(err, "Error: Failed to get platform\n"); - goto CLEANUP; - } + test_error_and_cleanup(err, CLEANUP, "Error: Failed to get platform\n"); err = clGetDeviceIDs(platform, CL_DEVICE_TYPE_GPU, 0, NULL, &num_devices); - if (CL_SUCCESS != err) - { - print_error(err, "clGetDeviceIDs failed in returning no. of devices\n"); - goto CLEANUP; - } + test_error_and_cleanup( + err, CLEANUP, "clGetDeviceIDs failed in returning no. of devices\n"); + devices = (cl_device_id *)malloc(num_devices * sizeof(cl_device_id)); if (NULL == devices) { - err = CL_OUT_OF_HOST_MEMORY; - print_error(err, "Unable to allocate memory for devices\n"); - goto CLEANUP; + test_fail_and_cleanup(err, CLEANUP, + "Unable to allocate memory for devices\n"); } err = clGetDeviceIDs(platform, CL_DEVICE_TYPE_GPU, num_devices, devices, NULL); - if (CL_SUCCESS != err) - { - print_error(err, "Failed to get deviceID.\n"); - goto CLEANUP; - } + test_error_and_cleanup(err, CLEANUP, "Failed to get deviceID.\n"); + contextProperties[1] = (cl_context_properties)platform; log_info("Assigned contextproperties for platform\n"); for (device_no = 0; device_no < num_devices; device_no++) @@ -1360,12 +1365,20 @@ int test_image_common(cl_device_id device_, cl_context context_, goto CLEANUP; } err = clGetDeviceInfo(devices[device_no], CL_DEVICE_UUID_KHR, - CL_UUID_SIZE_KHR, uuid, &extensionSize); - if (CL_SUCCESS != err) + CL_UUID_SIZE_KHR, uuid, NULL); + test_error_and_cleanup(err, CLEANUP, + "clGetDeviceInfo failed with error"); + + supportedSemaphoreTypes = + getSupportedInteropExternalSemaphoreHandleTypes(devices[device_no], + vkDevice); + + // If device does not support any semaphores, try the next one + if (supportedSemaphoreTypes.empty()) { - print_error(err, "clGetDeviceInfo failed with error"); - goto CLEANUP; + continue; } + err = memcmp(uuid, vkDevice.getPhysicalDevice().getUUID(), VK_UUID_SIZE); if (err == 0) @@ -1373,48 +1386,41 @@ int test_image_common(cl_device_id device_, cl_context context_, break; } } + + if (supportedSemaphoreTypes.empty()) + { + test_fail_and_cleanup( + err, CLEANUP, "No devices found that support OpenCL semaphores\n"); + } + if (device_no >= num_devices) { - err = EXIT_FAILURE; - print_error(err, - "OpenCL error:" - "No Vulkan-OpenCL Interop capable GPU found.\n"); - goto CLEANUP; + test_fail_and_cleanup(err, CLEANUP, + "OpenCL error:" + "No Vulkan-OpenCL Interop capable GPU found.\n"); } deviceId = devices[device_no]; err = setMaxImageDimensions(deviceId, max_width, max_height); - if (CL_SUCCESS != err) - { - print_error(err, "error setting max image dimensions"); - goto CLEANUP; - } + test_error_and_cleanup(err, CLEANUP, "error setting max image dimensions"); + log_info("Set max_width to %lu and max_height to %lu\n", max_width, max_height); context = clCreateContextFromType(contextProperties, CL_DEVICE_TYPE_GPU, NULL, NULL, &err); - if (CL_SUCCESS != err) - { - print_error(err, "error creating context"); - goto CLEANUP; - } + test_error_and_cleanup(err, CLEANUP, "error creating context"); + log_info("Successfully created context !!!\n"); cmd_queue1 = clCreateCommandQueue(context, devices[device_no], 0, &err); - if (CL_SUCCESS != err) - { - err = CL_INVALID_COMMAND_QUEUE; - print_error(err, "Error: Failed to create command queue!\n"); - goto CLEANUP; - } + test_error_and_cleanup(err, CLEANUP, + "Error: Failed to create command queue!\n"); + log_info("clCreateCommandQueue successfull \n"); cmd_queue2 = clCreateCommandQueue(context, devices[device_no], 0, &err); - if (CL_SUCCESS != err) - { - err = CL_INVALID_COMMAND_QUEUE; - print_error(err, "Error: Failed to create command queue!\n"); - goto CLEANUP; - } + test_error_and_cleanup(err, CLEANUP, + "Error: Failed to create command queue!\n"); + log_info("clCreateCommandQueue2 successful \n"); for (int i = 0; i < num_kernels; i++) @@ -1473,43 +1479,34 @@ int test_image_common(cl_device_id device_, cl_context context_, context, 1, &sourceTexts[k], &program_source_length, &err); err |= clBuildProgram(program[k], 0, NULL, NULL, NULL, NULL); } + test_error_and_cleanup(err, CLEANUP, "Error: Failed to build program"); - if (err != CL_SUCCESS) - { - print_error(err, "Error: Failed to build program"); - goto CLEANUP; - } // create the kernel kernel_float[i] = clCreateKernel(program[0], "image2DKernel", &err); - if (err != CL_SUCCESS) - { - print_error(err, "clCreateKernel failed"); - goto CLEANUP; - } + test_error_and_cleanup(err, CLEANUP, "clCreateKernel failed"); + kernel_signed[i] = clCreateKernel(program[1], "image2DKernel", &err); - if (err != CL_SUCCESS) + test_error_and_cleanup(err, CLEANUP, "clCreateKernel failed"); + + kernel_unsigned[i] = clCreateKernel(program[2], "image2DKernel", &err); + test_error_and_cleanup(err, CLEANUP, "clCreateKernel failed "); + } + for (VulkanExternalSemaphoreHandleType externalSemaphoreType : + supportedSemaphoreTypes) + { + if (numCQ == 2) { - print_error(err, "clCreateKernel failed"); - goto CLEANUP; + err = run_test_with_two_queue( + context, cmd_queue1, cmd_queue2, kernel_unsigned, kernel_signed, + kernel_float, vkDevice, externalSemaphoreType); } - kernel_unsigned[i] = clCreateKernel(program[2], "image2DKernel", &err); - if (err != CL_SUCCESS) + else { - print_error(err, "clCreateKernel failed "); - goto CLEANUP; + err = run_test_with_one_queue(context, cmd_queue1, kernel_unsigned, + kernel_signed, kernel_float, vkDevice, + externalSemaphoreType); } } - if (numCQ == 2) - { - err = run_test_with_two_queue(context, cmd_queue1, cmd_queue2, - kernel_unsigned, kernel_signed, - kernel_float, vkDevice); - } - else - { - err = run_test_with_one_queue(context, cmd_queue1, kernel_unsigned, - kernel_signed, kernel_float, vkDevice); - } CLEANUP: for (int i = 0; i < num_kernels; i++) { diff --git a/test_conformance/vulkan/vulkan_interop_common.hpp b/test_conformance/vulkan/vulkan_interop_common.hpp index a116240702..6f4d91555f 100644 --- a/test_conformance/vulkan/vulkan_interop_common.hpp +++ b/test_conformance/vulkan/vulkan_interop_common.hpp @@ -43,7 +43,5 @@ extern bool debug_trace; extern bool useSingleImageKernel; extern bool useDeviceLocal; extern bool disableNTHandleType; -// Enable offset for multiImport of vulkan device memory -extern bool enableOffset; #endif // _vulkan_interop_common_hpp_