From d058dfdeef5a56278cf1e5b2321abb3f41873fea Mon Sep 17 00:00:00 2001 From: Marcin Hajder Date: Tue, 7 Jan 2025 19:09:38 +0100 Subject: [PATCH] Corrected test_vulkan to use specific platform/device from harness (#2154) Fixes #1926 according to task description --- .../vulkan_wrapper/opencl_vulkan_wrapper.cpp | 5 +- .../common/vulkan_wrapper/vulkan_utility.cpp | 86 +- .../common/vulkan_wrapper/vulkan_utility.hpp | 11 +- .../common/vulkan_wrapper/vulkan_wrapper.hpp | 2 +- test_conformance/vulkan/CMakeLists.txt | 1 + test_conformance/vulkan/main.cpp | 232 +-- test_conformance/vulkan/procs.h | 33 + .../vulkan/test_vulkan_api_consistency.cpp | 949 +++++++------ ...st_vulkan_api_consistency_for_1dimages.cpp | 363 ++--- ...st_vulkan_api_consistency_for_3dimages.cpp | 367 ++--- .../vulkan/test_vulkan_interop_buffer.cpp | 1238 +++++++++-------- .../vulkan/test_vulkan_interop_image.cpp | 422 +++--- .../test_vulkan_platform_device_info.cpp | 337 +++-- test_conformance/vulkan/vulkan_test_base.h | 129 ++ 14 files changed, 2090 insertions(+), 2085 deletions(-) create mode 100644 test_conformance/vulkan/vulkan_test_base.h diff --git a/test_conformance/common/vulkan_wrapper/opencl_vulkan_wrapper.cpp b/test_conformance/common/vulkan_wrapper/opencl_vulkan_wrapper.cpp index b4330e92..ded1e709 100644 --- a/test_conformance/common/vulkan_wrapper/opencl_vulkan_wrapper.cpp +++ b/test_conformance/common/vulkan_wrapper/opencl_vulkan_wrapper.cpp @@ -464,7 +464,7 @@ getCLImageInfoFromVkImageInfo(const VkImageCreateInfo *VulkanImageCreateInfo, memcpy(img_fmt, &clImgFormat, sizeof(cl_image_format)); img_desc->image_type = getImageTypeFromVk(VulkanImageCreateInfo->imageType); - if (CL_INVALID_VALUE == img_desc->image_type) + if (CL_INVALID_VALUE == static_cast(img_desc->image_type)) { return CL_INVALID_VALUE; } @@ -503,6 +503,8 @@ cl_int check_external_memory_handle_type( errNum = clGetDeviceInfo(deviceID, CL_DEVICE_EXTERNAL_MEMORY_IMPORT_HANDLE_TYPES_KHR, 0, NULL, &handle_type_size); + test_error(errNum, "clGetDeviceInfo failed"); + handle_type = (cl_external_memory_handle_type_khr *)malloc(handle_type_size); @@ -539,6 +541,7 @@ cl_int check_external_semaphore_handle_type( errNum = clGetDeviceInfo(deviceID, queryParamName, 0, NULL, &handle_type_size); + test_error(errNum, "clGetDeviceInfo failed"); if (handle_type_size == 0) { diff --git a/test_conformance/common/vulkan_wrapper/vulkan_utility.cpp b/test_conformance/common/vulkan_wrapper/vulkan_utility.cpp index 1c433a71..e4796bc9 100644 --- a/test_conformance/common/vulkan_wrapper/vulkan_utility.cpp +++ b/test_conformance/common/vulkan_wrapper/vulkan_utility.cpp @@ -1,5 +1,5 @@ // -// Copyright (c) 2022 The Khronos Group Inc. +// Copyright (c) 2024 The Khronos Group Inc. // // Licensed under the Apache License, Version 2.0 (the "License"); // you may not use this file except in compliance with the License. @@ -40,13 +40,10 @@ const VulkanInstance &getVulkanInstance() const VulkanPhysicalDevice &getVulkanPhysicalDevice() { - size_t pdIdx; + size_t pdIdx = 0; cl_int errNum = 0; - cl_platform_id platform = NULL; + cl_platform_id platform = nullptr; cl_uchar uuid[CL_UUID_SIZE_KHR]; - cl_device_id *devices; - char *extensions = NULL; - size_t extensionSize = 0; cl_uint num_devices = 0; cl_uint device_no = 0; const size_t bufsize = BUFFERSIZE; @@ -69,14 +66,9 @@ const VulkanPhysicalDevice &getVulkanPhysicalDevice() throw std::runtime_error( "Error: clGetDeviceIDs failed in returning of devices\n"); } - devices = (cl_device_id *)malloc(num_devices * sizeof(cl_device_id)); - if (NULL == devices) - { - throw std::runtime_error( - "Error: Unable to allocate memory for devices\n"); - } - errNum = clGetDeviceIDs(platform, CL_DEVICE_TYPE_GPU, num_devices, devices, - NULL); + std::vector devices(num_devices); + errNum = clGetDeviceIDs(platform, CL_DEVICE_TYPE_GPU, num_devices, + devices.data(), NULL); if (CL_SUCCESS != errNum) { throw std::runtime_error("Error: Failed to get deviceID.\n"); @@ -84,34 +76,14 @@ const VulkanPhysicalDevice &getVulkanPhysicalDevice() bool is_selected = false; 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) - { - throw std::runtime_error("Error in clGetDeviceInfo for getting " - "device_extension size....\n"); - } - extensions = (char *)malloc(extensionSize); - if (NULL == extensions) - { - throw std::runtime_error( - "Unable to allocate memory for extensions\n"); - } - errNum = clGetDeviceInfo(devices[device_no], CL_DEVICE_EXTENSIONS, - extensionSize, extensions, NULL); - if (CL_SUCCESS != errNum) - { - throw std::runtime_error("Error: Error in clGetDeviceInfo for " - "getting device_extension\n"); - } errNum = clGetDeviceInfo(devices[device_no], CL_DEVICE_UUID_KHR, - CL_UUID_SIZE_KHR, uuid, &extensionSize); + CL_UUID_SIZE_KHR, uuid, nullptr); if (CL_SUCCESS != errNum) { throw std::runtime_error( "Error: clGetDeviceInfo failed with error\n"); } - free(extensions); + for (pdIdx = 0; pdIdx < physicalDeviceList.size(); pdIdx++) { if (!memcmp(&uuid, physicalDeviceList[pdIdx].getUUID(), @@ -139,10 +111,48 @@ const VulkanPhysicalDevice &getVulkanPhysicalDevice() return physicalDeviceList[pdIdx]; } -const VulkanQueueFamily &getVulkanQueueFamily(uint32_t queueFlags) +const VulkanPhysicalDevice & +getAssociatedVulkanPhysicalDevice(cl_device_id deviceId) +{ + size_t pdIdx; + cl_int errNum = 0; + cl_uchar uuid[CL_UUID_SIZE_KHR]; + const VulkanInstance &instance = getVulkanInstance(); + const VulkanPhysicalDeviceList &physicalDeviceList = + instance.getPhysicalDeviceList(); + + errNum = clGetDeviceInfo(deviceId, CL_DEVICE_UUID_KHR, CL_UUID_SIZE_KHR, + uuid, nullptr); + if (CL_SUCCESS != errNum) + { + throw std::runtime_error("Error: clGetDeviceInfo failed with error\n"); + } + for (pdIdx = 0; pdIdx < physicalDeviceList.size(); pdIdx++) + { + if (!memcmp(&uuid, physicalDeviceList[pdIdx].getUUID(), VK_UUID_SIZE)) + { + std::cout << "Selected physical device = " + << physicalDeviceList[pdIdx] << std::endl; + break; + } + } + + if ((pdIdx >= physicalDeviceList.size()) + || (physicalDeviceList[pdIdx] == (VkPhysicalDevice)VK_NULL_HANDLE)) + { + throw std::runtime_error("failed to find a suitable GPU!"); + } + std::cout << "Selected physical device is: " << physicalDeviceList[pdIdx] + << std::endl; + return physicalDeviceList[pdIdx]; +} + + +const VulkanQueueFamily & +getVulkanQueueFamily(const VulkanPhysicalDevice &physicalDevice, + uint32_t queueFlags) { size_t qfIdx; - const VulkanPhysicalDevice &physicalDevice = getVulkanPhysicalDevice(); const VulkanQueueFamilyList &queueFamilyList = physicalDevice.getQueueFamilyList(); diff --git a/test_conformance/common/vulkan_wrapper/vulkan_utility.hpp b/test_conformance/common/vulkan_wrapper/vulkan_utility.hpp index d2f4b7bf..486ad97c 100644 --- a/test_conformance/common/vulkan_wrapper/vulkan_utility.hpp +++ b/test_conformance/common/vulkan_wrapper/vulkan_utility.hpp @@ -1,5 +1,5 @@ // -// Copyright (c) 2022 The Khronos Group Inc. +// Copyright (c) 2024 The Khronos Group Inc. // // Licensed under the Apache License, Version 2.0 (the "License"); // you may not use this file except in compliance with the License. @@ -32,9 +32,12 @@ const VulkanInstance& getVulkanInstance(); const VulkanPhysicalDevice& getVulkanPhysicalDevice(); -const VulkanQueueFamily& -getVulkanQueueFamily(uint32_t queueFlags = VULKAN_QUEUE_FLAG_GRAPHICS - | VULKAN_QUEUE_FLAG_COMPUTE); +const VulkanPhysicalDevice& +getAssociatedVulkanPhysicalDevice(cl_device_id deviceId); +const VulkanQueueFamily& getVulkanQueueFamily( + const VulkanPhysicalDevice& physicalDevice = getVulkanPhysicalDevice(), + uint32_t queueFlags = VULKAN_QUEUE_FLAG_GRAPHICS + | VULKAN_QUEUE_FLAG_COMPUTE); const VulkanMemoryType& getVulkanMemoryType(const VulkanDevice& device, VulkanMemoryTypeProperty memoryTypeProperty); diff --git a/test_conformance/common/vulkan_wrapper/vulkan_wrapper.hpp b/test_conformance/common/vulkan_wrapper/vulkan_wrapper.hpp index b528f4aa..a536d140 100644 --- a/test_conformance/common/vulkan_wrapper/vulkan_wrapper.hpp +++ b/test_conformance/common/vulkan_wrapper/vulkan_wrapper.hpp @@ -145,7 +145,7 @@ public: virtual ~VulkanDevice(); const VulkanPhysicalDevice &getPhysicalDevice() const; VulkanQueue & - getQueue(const VulkanQueueFamily &queueFamily = getVulkanQueueFamily(), + getQueue(const VulkanQueueFamily &queueFamily /* = getVulkanQueueFamily()*/, uint32_t queueIndex = 0); operator VkDevice() const; }; diff --git a/test_conformance/vulkan/CMakeLists.txt b/test_conformance/vulkan/CMakeLists.txt index 7b9bcc18..85313c39 100644 --- a/test_conformance/vulkan/CMakeLists.txt +++ b/test_conformance/vulkan/CMakeLists.txt @@ -25,6 +25,7 @@ set (${MODULE_NAME}_SOURCES test_vulkan_api_consistency_for_1dimages.cpp test_vulkan_platform_device_info.cpp vulkan_interop_common.cpp + vulkan_test_base.h ) include_directories("../common/vulkan_wrapper") diff --git a/test_conformance/vulkan/main.cpp b/test_conformance/vulkan/main.cpp index 7b1cf01c..7be31b23 100644 --- a/test_conformance/vulkan/main.cpp +++ b/test_conformance/vulkan/main.cpp @@ -1,5 +1,5 @@ // -// Copyright (c) 2022 The Khronos Group Inc. +// Copyright (c) 2024 The Khronos Group Inc. // // Licensed under the Apache License, Version 2.0 (the "License"); // you may not use this file except in compliance with the License. @@ -30,121 +30,15 @@ #include #endif - #include "procs.h" #include "harness/testHarness.h" -#include "harness/parseParameters.h" -#include "harness/deviceInfo.h" #if !defined(_WIN32) #include #endif -#include -#include #define BUFFERSIZE 3000 -static void params_reset() -{ - numCQ = 1; - multiImport = false; - multiCtx = false; -} - -extern int test_buffer_common(cl_device_id device_, cl_context context_, - cl_command_queue queue_, int numElements_, - bool use_fence); -extern int test_image_common(cl_device_id device_, cl_context context_, - cl_command_queue queue_, int numElements_); - -int test_buffer_single_queue(cl_device_id device_, cl_context context_, - cl_command_queue queue_, int numElements_) -{ - params_reset(); - log_info("RUNNING TEST WITH ONE QUEUE...... \n\n"); - return test_buffer_common(device_, context_, queue_, numElements_, false); -} -int test_buffer_multiple_queue(cl_device_id device_, cl_context context_, - cl_command_queue queue_, int numElements_) -{ - params_reset(); - numCQ = 2; - log_info("RUNNING TEST WITH TWO QUEUE...... \n\n"); - return test_buffer_common(device_, context_, queue_, numElements_, false); -} -int test_buffer_multiImport_sameCtx(cl_device_id device_, cl_context context_, - cl_command_queue queue_, int numElements_) -{ - params_reset(); - multiImport = true; - log_info("RUNNING TEST WITH MULTIPLE DEVICE MEMORY IMPORT " - "IN SAME CONTEXT...... \n\n"); - return test_buffer_common(device_, context_, queue_, numElements_, false); -} -int test_buffer_multiImport_diffCtx(cl_device_id device_, cl_context context_, - cl_command_queue queue_, int numElements_) -{ - params_reset(); - multiImport = true; - multiCtx = true; - log_info("RUNNING TEST WITH MULTIPLE DEVICE MEMORY IMPORT " - "IN DIFFERENT CONTEXT...... \n\n"); - return test_buffer_common(device_, context_, queue_, numElements_, false); -} -int test_buffer_single_queue_fence(cl_device_id device_, cl_context context_, - cl_command_queue queue_, int numElements_) -{ - params_reset(); - log_info("RUNNING TEST WITH ONE QUEUE...... \n\n"); - return test_buffer_common(device_, context_, queue_, numElements_, true); -} -int test_buffer_multiple_queue_fence(cl_device_id device_, cl_context context_, - cl_command_queue queue_, int numElements_) -{ - params_reset(); - numCQ = 2; - log_info("RUNNING TEST WITH TWO QUEUE...... \n\n"); - return test_buffer_common(device_, context_, queue_, numElements_, true); -} -int test_buffer_multiImport_sameCtx_fence(cl_device_id device_, - cl_context context_, - cl_command_queue queue_, - int numElements_) -{ - params_reset(); - multiImport = true; - log_info("RUNNING TEST WITH MULTIPLE DEVICE MEMORY IMPORT " - "IN SAME CONTEXT...... \n\n"); - return test_buffer_common(device_, context_, queue_, numElements_, true); -} -int test_buffer_multiImport_diffCtx_fence(cl_device_id device_, - cl_context context_, - cl_command_queue queue_, - int numElements_) -{ - params_reset(); - multiImport = true; - multiCtx = true; - log_info("RUNNING TEST WITH MULTIPLE DEVICE MEMORY IMPORT " - "IN DIFFERENT CONTEXT...... \n\n"); - return test_buffer_common(device_, context_, queue_, numElements_, true); -} -int test_image_single_queue(cl_device_id device_, cl_context context_, - cl_command_queue queue_, int numElements_) -{ - params_reset(); - log_info("RUNNING TEST WITH ONE QUEUE...... \n\n"); - return test_image_common(device_, context_, queue_, numElements_); -} -int test_image_multiple_queue(cl_device_id device_, cl_context context_, - cl_command_queue queue_, int numElements_) -{ - params_reset(); - numCQ = 2; - log_info("RUNNING TEST WITH TWO QUEUE...... \n\n"); - return test_image_common(device_, context_, queue_, numElements_); -} - test_definition test_list[] = { ADD_TEST(buffer_single_queue), ADD_TEST(buffer_multiple_queue), ADD_TEST(buffer_multiImport_sameCtx), @@ -165,20 +59,6 @@ test_definition test_list[] = { ADD_TEST(buffer_single_queue), const int test_num = ARRAY_SIZE(test_list); -cl_device_type gDeviceType = CL_DEVICE_TYPE_DEFAULT; -char *choosen_platform_name = NULL; -cl_platform_id platform = NULL; -cl_int choosen_platform_index = -1; -char platform_name[1024] = ""; -cl_platform_id select_platform = NULL; -char *extensions = NULL; -size_t extensionSize = 0; -cl_uint num_devices = 0; -cl_uint device_no = 0; -cl_device_id *devices; -const size_t bufsize = BUFFERSIZE; -char buf[BUFFERSIZE]; -cl_uchar uuid[CL_UUID_SIZE_KHR]; unsigned int numCQ; bool multiImport; bool multiCtx; @@ -269,19 +149,7 @@ size_t parseParams(int argc, const char *argv[], const char **argList) int main(int argc, const char *argv[]) { - int errNum = 0; - test_start(); - params_reset(); - - if (!checkVkSupport()) - { - log_info("Vulkan supported GPU not found \n"); - log_info("TEST SKIPPED \n"); - return 0; - } - - VulkanDevice vkDevice; cl_device_type requestedDeviceType = CL_DEVICE_TYPE_GPU; char *force_cpu = getenv("CL_DEVICE_TYPE"); @@ -305,104 +173,10 @@ int main(int argc, const char *argv[]) log_info("Vulkan tests can only run on a GPU device.\n"); return 0; } - gDeviceType = CL_DEVICE_TYPE_GPU; const char **argList = (const char **)calloc(argc, sizeof(char *)); size_t argCount = parseParams(argc, argv, argList); if (argCount == 0) return 0; - // get the platform ID - errNum = clGetPlatformIDs(1, &platform, NULL); - if (errNum != CL_SUCCESS) - { - print_error(errNum, "Error: Failed to get platform\n"); - return errNum; - } - 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"); - return errNum; - } - devices = (cl_device_id *)malloc(num_devices * sizeof(cl_device_id)); - if (NULL == devices) - { - print_error(errNum, "Unable to allocate memory for devices\n"); - return CL_OUT_OF_HOST_MEMORY; - } - errNum = clGetDeviceIDs(platform, CL_DEVICE_TYPE_GPU, num_devices, devices, - NULL); - if (CL_SUCCESS != errNum) - { - print_error(errNum, "Failed to get deviceID.\n"); - return errNum; - } - 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) - { - log_error("Error in clGetDeviceInfo for getting " - "device_extension size....\n"); - return errNum; - } - extensions = (char *)malloc(extensionSize); - if (NULL == extensions) - { - log_error("Unable to allocate memory for extensions\n"); - return CL_OUT_OF_HOST_MEMORY; - } - errNum = - clGetDeviceInfo(devices[device_no], CL_DEVICE_EXTENSIONS, - extensionSize, extensions, NULL /*&extensionSize*/); - if (CL_SUCCESS != errNum) - { - print_error(errNum, - "Error in clGetDeviceInfo for getting " - "device_extension\n"); - return errNum; - } - errNum = clGetDeviceInfo(devices[device_no], CL_DEVICE_UUID_KHR, - CL_UUID_SIZE_KHR, uuid, &extensionSize); - if (CL_SUCCESS != errNum) - { - print_error(errNum, "clGetDeviceInfo failed with error\n "); - return errNum; - } - errNum = - memcmp(uuid, vkDevice.getPhysicalDevice().getUUID(), VK_UUID_SIZE); - if (errNum == 0) - { - break; - } - } - if (device_no >= num_devices) - { - fprintf(stderr, - "OpenCL error: " - "No Vulkan-OpenCL Interop capable GPU found.\n"); - } - if (!(is_extension_available(devices[device_no], "cl_khr_external_memory") - && is_extension_available(devices[device_no], - "cl_khr_external_semaphore"))) - { - log_info("Device does not support cl_khr_external_memory " - "or cl_khr_external_semaphore\n"); - log_info(" TEST SKIPPED\n"); - return CL_SUCCESS; - } - init_cl_vk_ext(platform, num_devices, devices); - - // Execute tests. - // Note: don't use the entire harness, because we have a different way of - // obtaining the device (via the context) - test_harness_config config{}; - config.forceNoContextCreation = true; - config.numElementsToUse = 1024; - config.queueProps = 0; - errNum = parseAndCallCommandLineTests(argCount, argList, devices[device_no], - test_num, test_list, config); - return errNum; -} \ No newline at end of file + return runTestHarness(argc, argv, test_num, test_list, false, 0); +} diff --git a/test_conformance/vulkan/procs.h b/test_conformance/vulkan/procs.h index d5465d7a..71fad68f 100644 --- a/test_conformance/vulkan/procs.h +++ b/test_conformance/vulkan/procs.h @@ -44,3 +44,36 @@ extern int test_platform_info(cl_device_id device, cl_context context, cl_command_queue queue, int num_elements); extern int test_device_info(cl_device_id device, cl_context context, cl_command_queue queue, int num_elements); +extern int test_buffer_single_queue(cl_device_id device_, cl_context context_, + cl_command_queue queue_, int numElements_); +extern int test_buffer_multiple_queue(cl_device_id device_, cl_context context_, + cl_command_queue queue_, + int numElements_); +extern int test_buffer_multiImport_sameCtx(cl_device_id device_, + cl_context context_, + cl_command_queue queue_, + int numElements_); +extern int test_buffer_multiImport_diffCtx(cl_device_id device_, + cl_context context_, + cl_command_queue queue_, + int numElements_); +extern int test_buffer_single_queue_fence(cl_device_id device_, + cl_context context_, + cl_command_queue queue_, + int numElements_); +extern int test_buffer_multiple_queue_fence(cl_device_id device_, + cl_context context_, + cl_command_queue queue_, + int numElements_); +extern int test_buffer_multiImport_sameCtx_fence(cl_device_id device_, + cl_context context_, + cl_command_queue queue_, + int numElements_); +extern int test_buffer_multiImport_diffCtx_fence(cl_device_id device_, + cl_context context_, + cl_command_queue queue_, + int numElements_); +extern int test_image_single_queue(cl_device_id device_, cl_context context_, + cl_command_queue queue_, int numElements_); +extern int test_image_multiple_queue(cl_device_id device_, cl_context context_, + cl_command_queue queue_, int numElements_); diff --git a/test_conformance/vulkan/test_vulkan_api_consistency.cpp b/test_conformance/vulkan/test_vulkan_api_consistency.cpp index 09e02981..b27a3c74 100644 --- a/test_conformance/vulkan/test_vulkan_api_consistency.cpp +++ b/test_conformance/vulkan/test_vulkan_api_consistency.cpp @@ -1,5 +1,5 @@ // -// Copyright (c) 2022 The Khronos Group Inc. +// Copyright (c) 2024 The Khronos Group Inc. // // Licensed under the Apache License, Version 2.0 (the "License"); // you may not use this file except in compliance with the License. @@ -33,502 +33,525 @@ #include "harness/typeWrappers.h" #include "harness/deviceInfo.h" -int test_consistency_external_buffer(cl_device_id deviceID, cl_context _context, - cl_command_queue _queue, int num_elements) +#include "vulkan_test_base.h" +#include "opencl_vulkan_wrapper.hpp" + +namespace { + +struct ConsistencyExternalBufferTest : public VulkanTestBase { - cl_int errNum; - VulkanDevice vkDevice; - // Context and command queue creation - cl_platform_id platform = NULL; - cl_context context = NULL; - cl_command_queue cmd_queue = NULL; + ConsistencyExternalBufferTest(cl_device_id device, cl_context context, + cl_command_queue queue, cl_int nelems) + : VulkanTestBase(device, context, queue, nelems) + {} - cl_context_properties contextProperties[] = { CL_CONTEXT_PLATFORM, 0, 0 }; - errNum = clGetPlatformIDs(1, &platform, NULL); - test_error(errNum, "Failed to get platform Id"); + cl_int Run() override + { - contextProperties[1] = (cl_context_properties)platform; - - context = clCreateContextFromType(contextProperties, CL_DEVICE_TYPE_GPU, - NULL, NULL, &errNum); - test_error(errNum, "Unable to create context with properties"); - - cmd_queue = clCreateCommandQueue(context, deviceID, 0, &errNum); - test_error(errNum, "Unable to create command queue"); - - uint32_t bufferSize = 32; - cl_device_id devList[] = { deviceID, NULL }; + cl_int errNum = CL_SUCCESS; + uint32_t bufferSize = 32; #ifdef _WIN32 - if (!is_extension_available(devList[0], "cl_khr_external_memory_win32")) - { - throw std::runtime_error("Device does not support " - "cl_khr_external_memory_win32 extension \n"); - } + if (!is_extension_available(device, "cl_khr_external_memory_win32")) + { + throw std::runtime_error( + "Device does not support " + "cl_khr_external_memory_win32 extension \n"); + } #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"); - } + if (!is_extension_available(device, "cl_khr_external_memory_opaque_fd")) + { + throw std::runtime_error( + "Device does not support " + "cl_khr_external_memory_opaque_fd extension \n"); + } #endif - VulkanExternalMemoryHandleType vkExternalMemoryHandleType = - getSupportedVulkanExternalMemoryHandleTypeList()[0]; + VulkanExternalMemoryHandleType vkExternalMemoryHandleType = + getSupportedVulkanExternalMemoryHandleTypeList()[0]; - VulkanBuffer vkDummyBuffer(vkDevice, 4 * 1024, vkExternalMemoryHandleType); - const VulkanMemoryTypeList& memoryTypeList = - vkDummyBuffer.getMemoryTypeList(); + VulkanBuffer vkDummyBuffer(*vkDevice, 4 * 1024, + vkExternalMemoryHandleType); + const VulkanMemoryTypeList& memoryTypeList = + vkDummyBuffer.getMemoryTypeList(); - VulkanBufferList vkBufferList(1, vkDevice, bufferSize, - vkExternalMemoryHandleType); - VulkanDeviceMemory* vkDeviceMem = - new VulkanDeviceMemory(vkDevice, vkBufferList[0], memoryTypeList[0], - vkExternalMemoryHandleType); + VulkanBufferList vkBufferList(1, *vkDevice, bufferSize, + vkExternalMemoryHandleType); + VulkanDeviceMemory* vkDeviceMem = new VulkanDeviceMemory( + *vkDevice, vkBufferList[0], memoryTypeList[0], + vkExternalMemoryHandleType); - vkDeviceMem->bindBuffer(vkBufferList[0], 0); + vkDeviceMem->bindBuffer(vkBufferList[0], 0); - void* handle = NULL; - int fd; + void* handle = NULL; + int fd; - std::vector extMemProperties{ - (cl_mem_properties)CL_MEM_DEVICE_HANDLE_LIST_KHR, - (cl_mem_properties)devList[0], - (cl_mem_properties)CL_MEM_DEVICE_HANDLE_LIST_END_KHR, - }; - cl_external_memory_handle_type_khr type; - switch (vkExternalMemoryHandleType) - { + std::vector extMemProperties{ + (cl_mem_properties)CL_MEM_DEVICE_HANDLE_LIST_KHR, + (cl_mem_properties)device, + (cl_mem_properties)CL_MEM_DEVICE_HANDLE_LIST_END_KHR, + }; + cl_external_memory_handle_type_khr type; + switch (vkExternalMemoryHandleType) + { #ifdef _WIN32 - case VULKAN_EXTERNAL_MEMORY_HANDLE_TYPE_OPAQUE_WIN32_NT: - handle = vkDeviceMem->getHandle(vkExternalMemoryHandleType); - type = CL_EXTERNAL_MEMORY_HANDLE_OPAQUE_WIN32_KHR; - errNum = check_external_memory_handle_type(devList[0], type); - extMemProperties.push_back((cl_mem_properties)type); - extMemProperties.push_back((cl_mem_properties)handle); - break; - case VULKAN_EXTERNAL_MEMORY_HANDLE_TYPE_OPAQUE_WIN32_KMT: - handle = vkDeviceMem->getHandle(vkExternalMemoryHandleType); - type = CL_EXTERNAL_MEMORY_HANDLE_OPAQUE_WIN32_KMT_KHR; - errNum = check_external_memory_handle_type(devList[0], type); - extMemProperties.push_back((cl_mem_properties)type); - extMemProperties.push_back((cl_mem_properties)handle); - break; + case VULKAN_EXTERNAL_MEMORY_HANDLE_TYPE_OPAQUE_WIN32_NT: + handle = vkDeviceMem->getHandle(vkExternalMemoryHandleType); + type = CL_EXTERNAL_MEMORY_HANDLE_OPAQUE_WIN32_KHR; + errNum = check_external_memory_handle_type(device, type); + extMemProperties.push_back((cl_mem_properties)type); + extMemProperties.push_back((cl_mem_properties)handle); + break; + case VULKAN_EXTERNAL_MEMORY_HANDLE_TYPE_OPAQUE_WIN32_KMT: + handle = vkDeviceMem->getHandle(vkExternalMemoryHandleType); + type = CL_EXTERNAL_MEMORY_HANDLE_OPAQUE_WIN32_KMT_KHR; + errNum = check_external_memory_handle_type(device, type); + extMemProperties.push_back((cl_mem_properties)type); + extMemProperties.push_back((cl_mem_properties)handle); + break; #else - case VULKAN_EXTERNAL_MEMORY_HANDLE_TYPE_OPAQUE_FD: - fd = (int)vkDeviceMem->getHandle(vkExternalMemoryHandleType); - type = CL_EXTERNAL_MEMORY_HANDLE_OPAQUE_FD_KHR; - errNum = check_external_memory_handle_type(devList[0], type); - extMemProperties.push_back((cl_mem_properties)type); - extMemProperties.push_back((cl_mem_properties)fd); - break; + case VULKAN_EXTERNAL_MEMORY_HANDLE_TYPE_OPAQUE_FD: + fd = (int)vkDeviceMem->getHandle(vkExternalMemoryHandleType); + type = CL_EXTERNAL_MEMORY_HANDLE_OPAQUE_FD_KHR; + errNum = check_external_memory_handle_type(device, type); + extMemProperties.push_back((cl_mem_properties)type); + extMemProperties.push_back((cl_mem_properties)fd); + break; #endif - default: - errNum = TEST_FAIL; - log_error("Unsupported external memory handle type \n"); - break; - } - if (errNum != CL_SUCCESS) - { - log_error("Checks failed for " - "CL_DEVICE_EXTERNAL_MEMORY_IMPORT_HANDLE_TYPES_KHR\n"); - return TEST_FAIL; - } - extMemProperties.push_back(0); + default: + errNum = TEST_FAIL; + log_error("Unsupported external memory handle type \n"); + break; + } + if (errNum != CL_SUCCESS) + { + log_error("Checks failed for " + "CL_DEVICE_EXTERNAL_MEMORY_IMPORT_HANDLE_TYPES_KHR\n"); + return TEST_FAIL; + } + extMemProperties.push_back(0); - clMemWrapper buffer; + clMemWrapper buffer; - // Passing NULL properties and a valid extMem_desc size - buffer = clCreateBufferWithProperties(context, NULL, 1, bufferSize, NULL, - &errNum); - test_error(errNum, "Unable to create buffer with NULL properties"); + // Passing NULL properties and a valid extMem_desc size + buffer = clCreateBufferWithProperties(context, NULL, 1, bufferSize, + NULL, &errNum); + test_error(errNum, "Unable to create buffer with NULL properties"); - buffer.reset(); + buffer.reset(); - // Passing valid extMemProperties and buffersize - buffer = clCreateBufferWithProperties(context, extMemProperties.data(), 1, - bufferSize, NULL, &errNum); - test_error(errNum, "Unable to create buffer with Properties"); + // Passing valid extMemProperties and buffersize + buffer = clCreateBufferWithProperties(context, extMemProperties.data(), + 1, bufferSize, NULL, &errNum); + test_error(errNum, "Unable to create buffer with Properties"); - buffer.reset(); + buffer.reset(); - // Not passing external memory handle - std::vector extMemProperties2{ + // Not passing external memory handle + std::vector extMemProperties2{ #ifdef _WIN32 - (cl_mem_properties)type, - NULL, // Passing NULL handle + (cl_mem_properties)type, + NULL, // Passing NULL handle #else - (cl_mem_properties)type, - (cl_mem_properties)-64, // Passing random invalid fd + (cl_mem_properties)type, + (cl_mem_properties)-64, // Passing random invalid fd #endif - (cl_mem_properties)CL_MEM_DEVICE_HANDLE_LIST_KHR, - (cl_mem_properties)devList[0], - (cl_mem_properties)CL_MEM_DEVICE_HANDLE_LIST_END_KHR, - 0 - }; - buffer = clCreateBufferWithProperties(context, extMemProperties2.data(), 1, - bufferSize, NULL, &errNum); - test_failure_error(errNum, CL_INVALID_VALUE, - "Should return CL_INVALID_VALUE "); + (cl_mem_properties)CL_MEM_DEVICE_HANDLE_LIST_KHR, + (cl_mem_properties)device, + (cl_mem_properties)CL_MEM_DEVICE_HANDLE_LIST_END_KHR, + 0 + }; + buffer = clCreateBufferWithProperties(context, extMemProperties2.data(), + 1, bufferSize, NULL, &errNum); + test_failure_error(errNum, CL_INVALID_VALUE, + "Should return CL_INVALID_VALUE "); - buffer.reset(); + buffer.reset(); - // Passing extMem_desc size = 0 but valid memProperties, CL_INVALID_SIZE - // should be returned. - buffer = clCreateBufferWithProperties(context, extMemProperties.data(), 1, - 0, NULL, &errNum); - test_failure_error(errNum, CL_INVALID_BUFFER_SIZE, - "Should return CL_INVALID_BUFFER_SIZE"); + // Passing extMem_desc size = 0 but valid memProperties, CL_INVALID_SIZE + // should be returned. + buffer = clCreateBufferWithProperties(context, extMemProperties.data(), + 1, 0, NULL, &errNum); + test_failure_error(errNum, CL_INVALID_BUFFER_SIZE, + "Should return CL_INVALID_BUFFER_SIZE"); - return TEST_PASS; + return TEST_PASS; + } +}; + +struct ConsistencyExternalImageTest : public VulkanTestBase +{ + ConsistencyExternalImageTest(cl_device_id device, cl_context context, + cl_command_queue queue, cl_int nelems) + : VulkanTestBase(device, context, queue, nelems) + {} + + cl_int Run() override + { + cl_int errNum = CL_SUCCESS; + +#ifdef _WIN32 + if (!is_extension_available(device, "cl_khr_external_memory_win32")) + { + throw std::runtime_error( + "Device does not support" + "cl_khr_external_memory_win32 extension \n"); + } +#else + if (!is_extension_available(device, "cl_khr_external_memory_opaque_fd")) + { + test_fail( + "Device does not support cl_khr_external_memory_opaque_fd " + "extension \n"); + } +#endif + uint32_t width = 256; + uint32_t height = 16; + cl_image_desc image_desc; + memset(&image_desc, 0x0, sizeof(cl_image_desc)); + cl_image_format img_format = { 0 }; + + VulkanExternalMemoryHandleType vkExternalMemoryHandleType = + getSupportedVulkanExternalMemoryHandleTypeList()[0]; + + VulkanImageTiling vulkanImageTiling = + vkClExternalMemoryHandleTilingAssumption( + device, vkExternalMemoryHandleType, &errNum); + ASSERT_SUCCESS(errNum, "Failed to query OpenCL tiling mode"); + + VulkanImage2D vkImage2D = VulkanImage2D( + *vkDevice, VULKAN_FORMAT_R8G8B8A8_UNORM, width, height, + vulkanImageTiling, 1, vkExternalMemoryHandleType); + + const VulkanMemoryTypeList& memoryTypeList = + vkImage2D.getMemoryTypeList(); + uint64_t totalImageMemSize = vkImage2D.getSize(); + + log_info("Memory type index: %u\n", (uint32_t)memoryTypeList[0]); + log_info("Memory type property: %d\n", + memoryTypeList[0].getMemoryTypeProperty()); + log_info("Image size : %ld\n", totalImageMemSize); + + VulkanDeviceMemory* vkDeviceMem = + new VulkanDeviceMemory(*vkDevice, vkImage2D, memoryTypeList[0], + vkExternalMemoryHandleType); + vkDeviceMem->bindImage(vkImage2D, 0); + + void* handle = NULL; + int fd; + std::vector extMemProperties{ + (cl_mem_properties)CL_MEM_DEVICE_HANDLE_LIST_KHR, + (cl_mem_properties)device, + (cl_mem_properties)CL_MEM_DEVICE_HANDLE_LIST_END_KHR, + }; + switch (vkExternalMemoryHandleType) + { +#ifdef _WIN32 + case VULKAN_EXTERNAL_MEMORY_HANDLE_TYPE_OPAQUE_WIN32_NT: + handle = vkDeviceMem->getHandle(vkExternalMemoryHandleType); + errNum = check_external_memory_handle_type( + device, CL_EXTERNAL_MEMORY_HANDLE_OPAQUE_WIN32_KHR); + extMemProperties.push_back( + (cl_mem_properties) + CL_EXTERNAL_MEMORY_HANDLE_OPAQUE_WIN32_KHR); + extMemProperties.push_back((cl_mem_properties)handle); + break; + case VULKAN_EXTERNAL_MEMORY_HANDLE_TYPE_OPAQUE_WIN32_KMT: + handle = vkDeviceMem->getHandle(vkExternalMemoryHandleType); + errNum = check_external_memory_handle_type( + device, CL_EXTERNAL_MEMORY_HANDLE_OPAQUE_WIN32_KMT_KHR); + extMemProperties.push_back( + (cl_mem_properties) + CL_EXTERNAL_MEMORY_HANDLE_OPAQUE_WIN32_KMT_KHR); + extMemProperties.push_back((cl_mem_properties)handle); + break; +#else + case VULKAN_EXTERNAL_MEMORY_HANDLE_TYPE_OPAQUE_FD: + fd = (int)vkDeviceMem->getHandle(vkExternalMemoryHandleType); + errNum = check_external_memory_handle_type( + device, CL_EXTERNAL_MEMORY_HANDLE_OPAQUE_FD_KHR); + extMemProperties.push_back( + (cl_mem_properties)CL_EXTERNAL_MEMORY_HANDLE_OPAQUE_FD_KHR); + extMemProperties.push_back((cl_mem_properties)fd); + break; +#endif + default: + errNum = TEST_FAIL; + log_error("Unsupported external memory handle type \n"); + break; + } + if (errNum != CL_SUCCESS) + { + log_error("Checks failed for " + "CL_DEVICE_EXTERNAL_MEMORY_IMPORT_HANDLE_TYPES_KHR\n"); + return TEST_FAIL; + } + extMemProperties.push_back(0); + + const VkImageCreateInfo VulkanImageCreateInfo = + vkImage2D.getVkImageCreateInfo(); + + errNum = getCLImageInfoFromVkImageInfo(&VulkanImageCreateInfo, + totalImageMemSize, &img_format, + &image_desc); + if (errNum != CL_SUCCESS) + { + log_error("getCLImageInfoFromVkImageInfo failed!!!"); + return TEST_FAIL; + } + + clMemWrapper image; + + // Pass valid properties, image_desc and image_format + image = clCreateImageWithProperties( + context, extMemProperties.data(), CL_MEM_READ_WRITE, &img_format, + &image_desc, NULL /* host_ptr */, &errNum); + test_error(errNum, "Unable to create Image with Properties"); + image.reset(); + + // Passing image_format as NULL + image = clCreateImageWithProperties(context, extMemProperties.data(), + CL_MEM_READ_WRITE, NULL, + &image_desc, NULL, &errNum); + test_failure_error(errNum, CL_INVALID_IMAGE_FORMAT_DESCRIPTOR, + "Image creation must fail with " + "CL_INVALID_IMAGE_FORMAT_DESCRIPTOR" + "when image desc passed as NULL"); + + image.reset(); + + // Passing image_desc as NULL + image = clCreateImageWithProperties(context, extMemProperties.data(), + CL_MEM_READ_WRITE, &img_format, + NULL, NULL, &errNum); + test_failure_error(errNum, CL_INVALID_IMAGE_DESCRIPTOR, + "Image creation must fail with " + "CL_INVALID_IMAGE_DESCRIPTOR " + "when image desc passed as NULL"); + image.reset(); + + return TEST_PASS; + } +}; + +struct ConsistencyExternalSemaphoreTest : public VulkanTestBase +{ + ConsistencyExternalSemaphoreTest(cl_device_id device, cl_context context, + cl_command_queue queue, cl_int nelems) + : VulkanTestBase(device, context, queue, nelems) + {} + + cl_int Run() override + { + cl_int errNum = CL_SUCCESS; + +#ifdef _WIN32 + if (!is_extension_available(device, "cl_khr_external_memory_win32")) + { + throw std::runtime_error( + "Device does not support" + "cl_khr_external_memory_win32 extension \n"); + } +#else + if (!is_extension_available(device, "cl_khr_external_memory_opaque_fd")) + { + test_fail( + "Device does not support cl_khr_external_memory_opaque_fd " + "extension \n"); + } +#endif + + std::vector + supportedExternalSemaphores = + getSupportedInteropExternalSemaphoreHandleTypes(device, + *vkDevice); + + if (supportedExternalSemaphores.empty()) + { + test_fail("No supported external semaphore types found\n"); + } + + 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(semaphoreHandleType); + handle2 = vkCl2Vksemaphore.getHandle(semaphoreHandleType); + errNum = check_external_semaphore_handle_type( + device, 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( + device, 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: + fd1 = (int)vkVk2Clsemaphore.getHandle(semaphoreHandleType); + fd2 = (int)vkCl2Vksemaphore.getHandle(semaphoreHandleType); + errNum = check_external_semaphore_handle_type( + device, 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( + device, 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_SEMAPHORE_DEVICE_HANDLE_LIST_KHR); + sema_props1.push_back((cl_semaphore_properties_khr)device); + sema_props1.push_back((cl_semaphore_properties_khr) + CL_SEMAPHORE_DEVICE_HANDLE_LIST_END_KHR); + sema_props2.push_back((cl_semaphore_properties_khr) + CL_SEMAPHORE_DEVICE_HANDLE_LIST_KHR); + sema_props2.push_back((cl_semaphore_properties_khr)device); + sema_props2.push_back((cl_semaphore_properties_khr) + CL_SEMAPHORE_DEVICE_HANDLE_LIST_END_KHR); + sema_props1.push_back(0); + sema_props2.push_back(0); + + // Pass NULL properties + clCreateSemaphoreWithPropertiesKHRptr(context, NULL, &errNum); + test_failure_error( + errNum, CL_INVALID_VALUE, + "Semaphore creation must fail with CL_INVALID_VALUE " + " when properties are passed as NULL"); + + // Pass invalid semaphore object to wait + errNum = clEnqueueWaitSemaphoresKHRptr(queue, 1, NULL, NULL, 0, + NULL, NULL); + test_failure_error( + errNum, CL_INVALID_VALUE, + "clEnqueueWaitSemaphoresKHR fails with CL_INVALID_VALUE " + "when invalid semaphore object is passed"); + + // Pass invalid semaphore object to signal + errNum = clEnqueueSignalSemaphoresKHRptr(queue, 1, NULL, NULL, 0, + NULL, NULL); + test_failure_error( + errNum, CL_INVALID_VALUE, + "clEnqueueSignalSemaphoresKHR fails with CL_INVALID_VALUE" + "when invalid semaphore object is passed"); + + // Create two semaphore objects + clVk2Clsemaphore = clCreateSemaphoreWithPropertiesKHRptr( + context, sema_props1.data(), &errNum); + test_error( + errNum, + "Unable to create semaphore with valid semaphore properties"); + + clCl2Vksemaphore = clCreateSemaphoreWithPropertiesKHRptr( + context, sema_props2.data(), &errNum); + test_error( + errNum, + "Unable to create semaphore with valid semaphore properties"); + + // Pass invalid object to release call + errNum = clReleaseSemaphoreKHRptr(NULL); + test_failure_error(errNum, CL_INVALID_SEMAPHORE_KHR, + "clReleaseSemaphoreKHRptr fails with " + "CL_INVALID_SEMAPHORE_KHR when NULL semaphore " + "object is passed"); + + // Release both semaphore objects + errNum = clReleaseSemaphoreKHRptr(clVk2Clsemaphore); + test_error(errNum, "clReleaseSemaphoreKHRptr failed"); + + errNum = clReleaseSemaphoreKHRptr(clCl2Vksemaphore); + test_error(errNum, "clReleaseSemaphoreKHRptr failed"); + } + + return TEST_PASS; + } +}; + +} // anonymous namespace + +int test_consistency_external_buffer(cl_device_id deviceID, cl_context context, + cl_command_queue defaultQueue, + int num_elements) +{ + return MakeAndRunTest( + deviceID, context, defaultQueue, num_elements); } -int test_consistency_external_image(cl_device_id deviceID, cl_context _context, - cl_command_queue _queue, int num_elements) +int test_consistency_external_image(cl_device_id deviceID, cl_context context, + cl_command_queue defaultQueue, + int num_elements) { - cl_int errNum; - VulkanDevice vkDevice; - - // Context and command queue creation - cl_platform_id platform = NULL; - cl_context context = NULL; - cl_command_queue cmd_queue = NULL; - - cl_context_properties contextProperties[] = { CL_CONTEXT_PLATFORM, 0, 0 }; - errNum = clGetPlatformIDs(1, &platform, NULL); - test_error(errNum, "Failed to get platform id"); - - contextProperties[1] = (cl_context_properties)platform; - - context = clCreateContextFromType(contextProperties, CL_DEVICE_TYPE_GPU, - NULL, NULL, &errNum); - test_error(errNum, "Unable to create context with properties"); - - cmd_queue = clCreateCommandQueue(context, deviceID, 0, &errNum); - test_error(errNum, "Unable to create command queue"); - - cl_device_id devList[] = { deviceID, NULL }; - -#ifdef _WIN32 - if (!is_extension_available(devList[0], "cl_khr_external_memory_win32")) - { - throw std::runtime_error("Device does not support" - "cl_khr_external_memory_win32 extension \n"); - } -#else - if (!is_extension_available(devList[0], "cl_khr_external_memory_opaque_fd")) - { - test_fail("Device does not support cl_khr_external_memory_opaque_fd " - "extension \n"); - } -#endif - uint32_t width = 256; - uint32_t height = 16; - cl_image_desc image_desc; - memset(&image_desc, 0x0, sizeof(cl_image_desc)); - cl_image_format img_format = { 0 }; - - VulkanExternalMemoryHandleType vkExternalMemoryHandleType = - getSupportedVulkanExternalMemoryHandleTypeList()[0]; - - VulkanImageTiling vulkanImageTiling = - vkClExternalMemoryHandleTilingAssumption( - deviceID, vkExternalMemoryHandleType, &errNum); - ASSERT_SUCCESS(errNum, "Failed to query OpenCL tiling mode"); - - VulkanImage2D vkImage2D = - VulkanImage2D(vkDevice, VULKAN_FORMAT_R8G8B8A8_UNORM, width, height, - vulkanImageTiling, 1, vkExternalMemoryHandleType); - - const VulkanMemoryTypeList& memoryTypeList = vkImage2D.getMemoryTypeList(); - uint64_t totalImageMemSize = vkImage2D.getSize(); - - log_info("Memory type index: %lu\n", (uint32_t)memoryTypeList[0]); - log_info("Memory type property: %d\n", - memoryTypeList[0].getMemoryTypeProperty()); - log_info("Image size : %d\n", totalImageMemSize); - - VulkanDeviceMemory* vkDeviceMem = new VulkanDeviceMemory( - vkDevice, vkImage2D, memoryTypeList[0], vkExternalMemoryHandleType); - vkDeviceMem->bindImage(vkImage2D, 0); - - void* handle = NULL; - int fd; - std::vector extMemProperties{ - (cl_mem_properties)CL_MEM_DEVICE_HANDLE_LIST_KHR, - (cl_mem_properties)devList[0], - (cl_mem_properties)CL_MEM_DEVICE_HANDLE_LIST_END_KHR, - }; - switch (vkExternalMemoryHandleType) - { -#ifdef _WIN32 - case VULKAN_EXTERNAL_MEMORY_HANDLE_TYPE_OPAQUE_WIN32_NT: - handle = vkDeviceMem->getHandle(vkExternalMemoryHandleType); - errNum = check_external_memory_handle_type( - devList[0], CL_EXTERNAL_MEMORY_HANDLE_OPAQUE_WIN32_KHR); - extMemProperties.push_back( - (cl_mem_properties)CL_EXTERNAL_MEMORY_HANDLE_OPAQUE_WIN32_KHR); - extMemProperties.push_back((cl_mem_properties)handle); - break; - case VULKAN_EXTERNAL_MEMORY_HANDLE_TYPE_OPAQUE_WIN32_KMT: - handle = vkDeviceMem->getHandle(vkExternalMemoryHandleType); - errNum = check_external_memory_handle_type( - devList[0], CL_EXTERNAL_MEMORY_HANDLE_OPAQUE_WIN32_KMT_KHR); - extMemProperties.push_back( - (cl_mem_properties) - CL_EXTERNAL_MEMORY_HANDLE_OPAQUE_WIN32_KMT_KHR); - extMemProperties.push_back((cl_mem_properties)handle); - break; -#else - case VULKAN_EXTERNAL_MEMORY_HANDLE_TYPE_OPAQUE_FD: - fd = (int)vkDeviceMem->getHandle(vkExternalMemoryHandleType); - errNum = check_external_memory_handle_type( - devList[0], CL_EXTERNAL_MEMORY_HANDLE_OPAQUE_FD_KHR); - extMemProperties.push_back( - (cl_mem_properties)CL_EXTERNAL_MEMORY_HANDLE_OPAQUE_FD_KHR); - extMemProperties.push_back((cl_mem_properties)fd); - break; -#endif - default: - errNum = TEST_FAIL; - log_error("Unsupported external memory handle type \n"); - break; - } - if (errNum != CL_SUCCESS) - { - log_error("Checks failed for " - "CL_DEVICE_EXTERNAL_MEMORY_IMPORT_HANDLE_TYPES_KHR\n"); - return TEST_FAIL; - } - extMemProperties.push_back(0); - - const VkImageCreateInfo VulkanImageCreateInfo = - vkImage2D.getVkImageCreateInfo(); - - errNum = getCLImageInfoFromVkImageInfo( - &VulkanImageCreateInfo, totalImageMemSize, &img_format, &image_desc); - if (errNum != CL_SUCCESS) - { - log_error("getCLImageInfoFromVkImageInfo failed!!!"); - return TEST_FAIL; - } - - clMemWrapper image; - - // Pass valid properties, image_desc and image_format - image = clCreateImageWithProperties( - context, extMemProperties.data(), CL_MEM_READ_WRITE, &img_format, - &image_desc, NULL /* host_ptr */, &errNum); - test_error(errNum, "Unable to create Image with Properties"); - image.reset(); - - // Passing image_format as NULL - image = clCreateImageWithProperties(context, extMemProperties.data(), - CL_MEM_READ_WRITE, NULL, &image_desc, - NULL, &errNum); - test_failure_error(errNum, CL_INVALID_IMAGE_FORMAT_DESCRIPTOR, - "Image creation must fail with " - "CL_INVALID_IMAGE_FORMAT_DESCRIPTOR" - "when image desc passed as NULL"); - - image.reset(); - - // Passing image_desc as NULL - image = clCreateImageWithProperties(context, extMemProperties.data(), - CL_MEM_READ_WRITE, &img_format, NULL, - NULL, &errNum); - test_failure_error(errNum, CL_INVALID_IMAGE_DESCRIPTOR, - "Image creation must fail with " - "CL_INVALID_IMAGE_DESCRIPTOR " - "when image desc passed as NULL"); - image.reset(); - - return TEST_PASS; + return MakeAndRunTest( + deviceID, context, defaultQueue, num_elements); } int test_consistency_external_semaphore(cl_device_id deviceID, - cl_context _context, - cl_command_queue _queue, + cl_context context, + cl_command_queue defaultQueue, int num_elements) { - cl_int errNum; - VulkanDevice vkDevice; - // Context and command queue creation - cl_platform_id platform = NULL; - cl_context context = NULL; - cl_command_queue cmd_queue = NULL; - - errNum = clGetPlatformIDs(1, &platform, NULL); - test_error(errNum, "Failed to get platform Id"); - - cl_context_properties contextProperties[] = { CL_CONTEXT_PLATFORM, 0, 0 }; - - contextProperties[1] = (cl_context_properties)platform; - - context = clCreateContextFromType(contextProperties, CL_DEVICE_TYPE_GPU, - NULL, NULL, &errNum); - test_error(errNum, "Unable to create context with properties"); - - cmd_queue = clCreateCommandQueue(context, deviceID, 0, &errNum); - test_error(errNum, "Unable to create command queue"); - - cl_device_id devList[] = { deviceID, NULL }; - - std::vector supportedExternalSemaphores = - getSupportedInteropExternalSemaphoreHandleTypes(devList[0], vkDevice); - - if (supportedExternalSemaphores.empty()) - { - test_fail("No supported external semaphore types found\n"); - } - - 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(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: - 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_SEMAPHORE_DEVICE_HANDLE_LIST_KHR); - sema_props1.push_back((cl_semaphore_properties_khr)devList[0]); - sema_props1.push_back( - (cl_semaphore_properties_khr)CL_SEMAPHORE_DEVICE_HANDLE_LIST_END_KHR); - sema_props2.push_back( - (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_SEMAPHORE_DEVICE_HANDLE_LIST_END_KHR); - sema_props1.push_back(0); - sema_props2.push_back(0); - - // Pass NULL properties - cl_semaphore_khr cl_ext_semaphore = - clCreateSemaphoreWithPropertiesKHRptr(context, NULL, &errNum); - test_failure_error(errNum, CL_INVALID_VALUE, - "Semaphore creation must fail with CL_INVALID_VALUE " - " when properties are passed as NULL"); - - - // Pass invalid semaphore object to wait - errNum = - clEnqueueWaitSemaphoresKHRptr(cmd_queue, 1, NULL, NULL, 0, NULL, NULL); - test_failure_error(errNum, CL_INVALID_VALUE, - "clEnqueueWaitSemaphoresKHR fails with CL_INVALID_VALUE " - "when invalid semaphore object is passed"); - - - // Pass invalid semaphore object to signal - errNum = clEnqueueSignalSemaphoresKHRptr(cmd_queue, 1, NULL, NULL, 0, NULL, - NULL); - test_failure_error( - errNum, CL_INVALID_VALUE, - "clEnqueueSignalSemaphoresKHR fails with CL_INVALID_VALUE" - "when invalid semaphore object is passed"); - - - // Create two semaphore objects - clVk2Clsemaphore = clCreateSemaphoreWithPropertiesKHRptr( - context, sema_props1.data(), &errNum); - test_error(errNum, - "Unable to create semaphore with valid semaphore properties"); - - clCl2Vksemaphore = clCreateSemaphoreWithPropertiesKHRptr( - context, sema_props2.data(), &errNum); - test_error(errNum, - "Unable to create semaphore with valid semaphore properties"); - - // Pass invalid object to release call - errNum = clReleaseSemaphoreKHRptr(NULL); - test_failure_error( - errNum, CL_INVALID_SEMAPHORE_KHR, - "clReleaseSemaphoreKHRptr fails with " - "CL_INVALID_SEMAPHORE_KHR when NULL semaphore object is passed"); - - // Release both semaphore objects - errNum = clReleaseSemaphoreKHRptr(clVk2Clsemaphore); - test_error(errNum, "clReleaseSemaphoreKHRptr failed"); - - errNum = clReleaseSemaphoreKHRptr(clCl2Vksemaphore); - test_error(errNum, "clReleaseSemaphoreKHRptr failed"); - } - - return TEST_PASS; + return MakeAndRunTest( + deviceID, context, defaultQueue, num_elements); } diff --git a/test_conformance/vulkan/test_vulkan_api_consistency_for_1dimages.cpp b/test_conformance/vulkan/test_vulkan_api_consistency_for_1dimages.cpp index aefdb414..799a73f0 100644 --- a/test_conformance/vulkan/test_vulkan_api_consistency_for_1dimages.cpp +++ b/test_conformance/vulkan/test_vulkan_api_consistency_for_1dimages.cpp @@ -1,3 +1,19 @@ +// +// Copyright (c) 2024 The Khronos Group Inc. +// +// Licensed under the Apache License, Version 2.0 (the "License"); +// you may not use this file except in compliance with the License. +// You may obtain a copy of the License at +// +// http://www.apache.org/licenses/LICENSE-2.0 +// +// Unless required by applicable law or agreed to in writing, software +// distributed under the License is distributed on an "AS IS" BASIS, +// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +// See the License for the specific language governing permissions and +// limitations under the License. +// + #include #include #include @@ -17,180 +33,181 @@ #include "harness/typeWrappers.h" #include "harness/deviceInfo.h" +#include "vulkan_test_base.h" +#include "opencl_vulkan_wrapper.hpp" + +namespace { + +struct ConsistencyExternalImage1DTest : public VulkanTestBase +{ + ConsistencyExternalImage1DTest(cl_device_id device, cl_context context, + cl_command_queue queue, cl_int nelems) + : VulkanTestBase(device, context, queue, nelems) + {} + + cl_int Run() override + { + cl_int errNum = CL_SUCCESS; + +#ifdef _WIN32 + if (!is_extension_available(device, "cl_khr_external_memory_win32")) + { + throw std::runtime_error( + "Device does not support" + "cl_khr_external_memory_win32 extension \n"); + } +#else + if (!is_extension_available(device, "cl_khr_external_memory_opaque_fd")) + { + throw std::runtime_error( + "Device does not support cl_khr_external_memory_opaque_fd " + "extension \n"); + } +#endif + uint32_t width = 256; + cl_image_desc image_desc; + memset(&image_desc, 0x0, sizeof(cl_image_desc)); + cl_image_format img_format = { 0 }; + + VulkanExternalMemoryHandleType vkExternalMemoryHandleType = + getSupportedVulkanExternalMemoryHandleTypeList()[0]; + + VulkanImageTiling vulkanImageTiling = + vkClExternalMemoryHandleTilingAssumption( + device, vkExternalMemoryHandleType, &errNum); + ASSERT_SUCCESS(errNum, "Failed to query OpenCL tiling mode"); + + VulkanImage1D vkImage1D = + VulkanImage1D(*vkDevice, VULKAN_FORMAT_R8G8B8A8_UNORM, width, + vulkanImageTiling, 1, vkExternalMemoryHandleType); + + const VulkanMemoryTypeList& memoryTypeList = + vkImage1D.getMemoryTypeList(); + uint64_t totalImageMemSize = vkImage1D.getSize(); + + log_info("Memory type index: %u\n", (uint32_t)memoryTypeList[0]); + log_info("Memory type property: %d\n", + memoryTypeList[0].getMemoryTypeProperty()); + log_info("Image size : %lu\n", totalImageMemSize); + + VulkanDeviceMemory* vkDeviceMem = + new VulkanDeviceMemory(*vkDevice, vkImage1D, memoryTypeList[0], + vkExternalMemoryHandleType); + vkDeviceMem->bindImage(vkImage1D, 0); + + void* handle = NULL; + int fd; + std::vector extMemProperties{ + (cl_mem_properties)CL_MEM_DEVICE_HANDLE_LIST_KHR, + (cl_mem_properties)device, + (cl_mem_properties)CL_MEM_DEVICE_HANDLE_LIST_END_KHR, + }; + switch (vkExternalMemoryHandleType) + { +#ifdef _WIN32 + case VULKAN_EXTERNAL_MEMORY_HANDLE_TYPE_OPAQUE_WIN32_NT: + handle = vkDeviceMem->getHandle(vkExternalMemoryHandleType); + errNum = check_external_memory_handle_type( + device, CL_EXTERNAL_MEMORY_HANDLE_OPAQUE_WIN32_KHR); + extMemProperties.push_back( + (cl_mem_properties) + CL_EXTERNAL_MEMORY_HANDLE_OPAQUE_WIN32_KHR); + extMemProperties.push_back((cl_mem_properties)handle); + break; + case VULKAN_EXTERNAL_MEMORY_HANDLE_TYPE_OPAQUE_WIN32_KMT: + handle = vkDeviceMem->getHandle(vkExternalMemoryHandleType); + errNum = check_external_memory_handle_type( + device, CL_EXTERNAL_MEMORY_HANDLE_OPAQUE_WIN32_KMT_KHR); + extMemProperties.push_back( + (cl_mem_properties) + CL_EXTERNAL_MEMORY_HANDLE_OPAQUE_WIN32_KMT_KHR); + extMemProperties.push_back((cl_mem_properties)handle); + break; +#else + case VULKAN_EXTERNAL_MEMORY_HANDLE_TYPE_OPAQUE_FD: + fd = (int)vkDeviceMem->getHandle(vkExternalMemoryHandleType); + errNum = check_external_memory_handle_type( + device, CL_EXTERNAL_MEMORY_HANDLE_OPAQUE_FD_KHR); + extMemProperties.push_back( + (cl_mem_properties)CL_EXTERNAL_MEMORY_HANDLE_OPAQUE_FD_KHR); + extMemProperties.push_back((cl_mem_properties)fd); + break; +#endif + default: + errNum = TEST_FAIL; + log_error("Unsupported external memory handle type \n"); + break; + } + if (errNum != CL_SUCCESS) + { + log_error("Checks failed for " + "CL_DEVICE_EXTERNAL_MEMORY_IMPORT_HANDLE_TYPES_KHR\n"); + return TEST_FAIL; + } + extMemProperties.push_back(0); + + const VkImageCreateInfo VulkanImageCreateInfo = + vkImage1D.getVkImageCreateInfo(); + + errNum = getCLImageInfoFromVkImageInfo(&VulkanImageCreateInfo, + totalImageMemSize, &img_format, + &image_desc); + if (errNum != CL_SUCCESS) + { + log_error("getCLImageInfoFromVkImageInfo failed!!!"); + return TEST_FAIL; + } + + clMemWrapper image; + + // Pass valid properties, image_desc and image_format + image = clCreateImageWithProperties( + context, extMemProperties.data(), CL_MEM_READ_WRITE, &img_format, + &image_desc, NULL /* host_ptr */, &errNum); + test_error(errNum, "Unable to create Image with Properties"); + 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, NULL, &errNum); + test_failure_error(errNum, CL_INVALID_IMAGE_FORMAT_DESCRIPTOR, + "Image creation must fail with " + "CL_INVALID_IMAGE_FORMAT_DESCRIPTOR" + "when image desc passed as NULL"); + + image.reset(); + + // Passing image_desc as NULL + image = clCreateImageWithProperties(context, extMemProperties.data(), + CL_MEM_READ_WRITE, &img_format, + NULL, NULL, &errNum); + test_failure_error(errNum, CL_INVALID_IMAGE_DESCRIPTOR, + "Image creation must fail with " + "CL_INVALID_IMAGE_DESCRIPTOR " + "when image desc passed as NULL"); + image.reset(); + + return TEST_PASS; + } +}; +} + int test_consistency_external_for_1dimage(cl_device_id deviceID, - cl_context _context, - cl_command_queue _queue, + cl_context context, + cl_command_queue defaultQueue, int num_elements) { - cl_int errNum; - VulkanDevice vkDevice; - - // Context and command queue creation - cl_platform_id platform = NULL; - cl_context context = NULL; - cl_command_queue cmd_queue = NULL; - - cl_context_properties contextProperties[] = { CL_CONTEXT_PLATFORM, 0, 0 }; - errNum = clGetPlatformIDs(1, &platform, NULL); - test_error(errNum, "Failed to get platform id"); - - contextProperties[1] = (cl_context_properties)platform; - - context = clCreateContextFromType(contextProperties, CL_DEVICE_TYPE_GPU, - NULL, NULL, &errNum); - test_error(errNum, "Unable to create context with properties"); - - cmd_queue = clCreateCommandQueue(context, deviceID, 0, &errNum); - test_error(errNum, "Unable to create command queue"); - - cl_device_id devList[] = { deviceID, NULL }; - -#ifdef _WIN32 - if (!is_extension_available(devList[0], "cl_khr_external_memory_win32")) - { - throw std::runtime_error("Device does not support" - "cl_khr_external_memory_win32 extension \n"); - } -#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"); - } -#endif - uint32_t width = 256; - cl_image_desc image_desc; - memset(&image_desc, 0x0, sizeof(cl_image_desc)); - cl_image_format img_format = { 0 }; - - VulkanExternalMemoryHandleType vkExternalMemoryHandleType = - getSupportedVulkanExternalMemoryHandleTypeList()[0]; - - VulkanImageTiling vulkanImageTiling = - vkClExternalMemoryHandleTilingAssumption( - deviceID, vkExternalMemoryHandleType, &errNum); - ASSERT_SUCCESS(errNum, "Failed to query OpenCL tiling mode"); - - VulkanImage1D vkImage1D = - VulkanImage1D(vkDevice, VULKAN_FORMAT_R8G8B8A8_UNORM, width, - vulkanImageTiling, 1, vkExternalMemoryHandleType); - - const VulkanMemoryTypeList& memoryTypeList = vkImage1D.getMemoryTypeList(); - uint64_t totalImageMemSize = vkImage1D.getSize(); - - log_info("Memory type index: %u\n", (uint32_t)memoryTypeList[0]); - log_info("Memory type property: %d\n", - memoryTypeList[0].getMemoryTypeProperty()); - log_info("Image size : %lu\n", totalImageMemSize); - - VulkanDeviceMemory* vkDeviceMem = new VulkanDeviceMemory( - vkDevice, vkImage1D, memoryTypeList[0], vkExternalMemoryHandleType); - vkDeviceMem->bindImage(vkImage1D, 0); - - void* handle = NULL; - int fd; - std::vector extMemProperties{ - (cl_mem_properties)CL_MEM_DEVICE_HANDLE_LIST_KHR, - (cl_mem_properties)devList[0], - (cl_mem_properties)CL_MEM_DEVICE_HANDLE_LIST_END_KHR, - }; - switch (vkExternalMemoryHandleType) - { -#ifdef _WIN32 - case VULKAN_EXTERNAL_MEMORY_HANDLE_TYPE_OPAQUE_WIN32_NT: - handle = vkDeviceMem->getHandle(vkExternalMemoryHandleType); - errNum = check_external_memory_handle_type( - devList[0], CL_EXTERNAL_MEMORY_HANDLE_OPAQUE_WIN32_KHR); - extMemProperties.push_back( - (cl_mem_properties)CL_EXTERNAL_MEMORY_HANDLE_OPAQUE_WIN32_KHR); - extMemProperties.push_back((cl_mem_properties)handle); - break; - case VULKAN_EXTERNAL_MEMORY_HANDLE_TYPE_OPAQUE_WIN32_KMT: - handle = vkDeviceMem->getHandle(vkExternalMemoryHandleType); - errNum = check_external_memory_handle_type( - devList[0], CL_EXTERNAL_MEMORY_HANDLE_OPAQUE_WIN32_KMT_KHR); - extMemProperties.push_back( - (cl_mem_properties) - CL_EXTERNAL_MEMORY_HANDLE_OPAQUE_WIN32_KMT_KHR); - extMemProperties.push_back((cl_mem_properties)handle); - break; -#else - case VULKAN_EXTERNAL_MEMORY_HANDLE_TYPE_OPAQUE_FD: - fd = (int)vkDeviceMem->getHandle(vkExternalMemoryHandleType); - errNum = check_external_memory_handle_type( - devList[0], CL_EXTERNAL_MEMORY_HANDLE_OPAQUE_FD_KHR); - extMemProperties.push_back( - (cl_mem_properties)CL_EXTERNAL_MEMORY_HANDLE_OPAQUE_FD_KHR); - extMemProperties.push_back((cl_mem_properties)fd); - break; -#endif - default: - errNum = TEST_FAIL; - log_error("Unsupported external memory handle type \n"); - break; - } - if (errNum != CL_SUCCESS) - { - log_error("Checks failed for " - "CL_DEVICE_EXTERNAL_MEMORY_IMPORT_HANDLE_TYPES_KHR\n"); - return TEST_FAIL; - } - extMemProperties.push_back(0); - - const VkImageCreateInfo VulkanImageCreateInfo = - vkImage1D.getVkImageCreateInfo(); - - errNum = getCLImageInfoFromVkImageInfo( - &VulkanImageCreateInfo, totalImageMemSize, &img_format, &image_desc); - if (errNum != CL_SUCCESS) - { - log_error("getCLImageInfoFromVkImageInfo failed!!!"); - return TEST_FAIL; - } - - clMemWrapper image; - - // Pass valid properties, image_desc and image_format - image = clCreateImageWithProperties( - context, extMemProperties.data(), CL_MEM_READ_WRITE, &img_format, - &image_desc, NULL /* host_ptr */, &errNum); - test_error(errNum, "Unable to create Image with Properties"); - 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, - NULL, &errNum); - test_failure_error(errNum, CL_INVALID_IMAGE_FORMAT_DESCRIPTOR, - "Image creation must fail with " - "CL_INVALID_IMAGE_FORMAT_DESCRIPTOR" - "when image desc passed as NULL"); - - image.reset(); - - // Passing image_desc as NULL - image = clCreateImageWithProperties(context, extMemProperties.data(), - CL_MEM_READ_WRITE, &img_format, NULL, - NULL, &errNum); - test_failure_error(errNum, CL_INVALID_IMAGE_DESCRIPTOR, - "Image creation must fail with " - "CL_INVALID_IMAGE_DESCRIPTOR " - "when image desc passed as NULL"); - image.reset(); - - if (cmd_queue) clReleaseCommandQueue(cmd_queue); - if (context) clReleaseContext(context); - - return TEST_PASS; + return MakeAndRunTest( + deviceID, context, defaultQueue, num_elements); } diff --git a/test_conformance/vulkan/test_vulkan_api_consistency_for_3dimages.cpp b/test_conformance/vulkan/test_vulkan_api_consistency_for_3dimages.cpp index a71fb945..b30f3747 100644 --- a/test_conformance/vulkan/test_vulkan_api_consistency_for_3dimages.cpp +++ b/test_conformance/vulkan/test_vulkan_api_consistency_for_3dimages.cpp @@ -1,3 +1,19 @@ +// +// Copyright (c) 2024 The Khronos Group Inc. +// +// Licensed under the Apache License, Version 2.0 (the "License"); +// you may not use this file except in compliance with the License. +// You may obtain a copy of the License at +// +// http://www.apache.org/licenses/LICENSE-2.0 +// +// Unless required by applicable law or agreed to in writing, software +// distributed under the License is distributed on an "AS IS" BASIS, +// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +// See the License for the specific language governing permissions and +// limitations under the License. +// + #include #include #include @@ -18,183 +34,184 @@ #include "harness/deviceInfo.h" #include +#include "vulkan_test_base.h" +#include "opencl_vulkan_wrapper.hpp" + +namespace { + +struct ConsistencyExternalImage3DTest : public VulkanTestBase +{ + ConsistencyExternalImage3DTest(cl_device_id device, cl_context context, + cl_command_queue queue, cl_int nelems) + : VulkanTestBase(device, context, queue, nelems) + {} + + cl_int Run() override + { + cl_int errNum; + +#ifdef _WIN32 + if (!is_extension_available(device, "cl_khr_external_memory_win32")) + { + throw std::runtime_error( + "Device does not support" + "cl_khr_external_memory_win32 extension \n"); + } +#else + if (!is_extension_available(device, "cl_khr_external_memory_opaque_fd")) + { + throw std::runtime_error( + "Device does not support cl_khr_external_memory_opaque_fd " + "extension \n"); + } +#endif + uint32_t width = 256; + uint32_t height = 16; + uint32_t depth = 10; + cl_image_desc image_desc; + memset(&image_desc, 0x0, sizeof(cl_image_desc)); + cl_image_format img_format = { 0 }; + + VulkanExternalMemoryHandleType vkExternalMemoryHandleType = + getSupportedVulkanExternalMemoryHandleTypeList()[0]; + + VulkanImageTiling vulkanImageTiling = + vkClExternalMemoryHandleTilingAssumption( + device, vkExternalMemoryHandleType, &errNum); + ASSERT_SUCCESS(errNum, "Failed to query OpenCL tiling mode"); + + VulkanImage3D vkImage3D = VulkanImage3D( + *vkDevice, VULKAN_FORMAT_R8G8B8A8_UNORM, width, height, depth, + vulkanImageTiling, 1, vkExternalMemoryHandleType); + + const VulkanMemoryTypeList& memoryTypeList = + vkImage3D.getMemoryTypeList(); + uint64_t totalImageMemSize = vkImage3D.getSize(); + + log_info("Memory type index: %u\n", (uint32_t)memoryTypeList[0]); + log_info("Memory type property: %d\n", + memoryTypeList[0].getMemoryTypeProperty()); + log_info("Image size : %lu\n", totalImageMemSize); + + VulkanDeviceMemory* vkDeviceMem = + new VulkanDeviceMemory(*vkDevice, vkImage3D, memoryTypeList[0], + vkExternalMemoryHandleType); + vkDeviceMem->bindImage(vkImage3D, 0); + + void* handle = NULL; + int fd; + std::vector extMemProperties{ + (cl_mem_properties)CL_MEM_DEVICE_HANDLE_LIST_KHR, + (cl_mem_properties)device, + (cl_mem_properties)CL_MEM_DEVICE_HANDLE_LIST_END_KHR, + }; + switch (vkExternalMemoryHandleType) + { +#ifdef _WIN32 + case VULKAN_EXTERNAL_MEMORY_HANDLE_TYPE_OPAQUE_WIN32_NT: + handle = vkDeviceMem->getHandle(vkExternalMemoryHandleType); + errNum = check_external_memory_handle_type( + device, CL_EXTERNAL_MEMORY_HANDLE_OPAQUE_WIN32_KHR); + extMemProperties.push_back( + (cl_mem_properties) + CL_EXTERNAL_MEMORY_HANDLE_OPAQUE_WIN32_KHR); + extMemProperties.push_back((cl_mem_properties)handle); + break; + case VULKAN_EXTERNAL_MEMORY_HANDLE_TYPE_OPAQUE_WIN32_KMT: + handle = vkDeviceMem->getHandle(vkExternalMemoryHandleType); + errNum = check_external_memory_handle_type( + device, CL_EXTERNAL_MEMORY_HANDLE_OPAQUE_WIN32_KMT_KHR); + extMemProperties.push_back( + (cl_mem_properties) + CL_EXTERNAL_MEMORY_HANDLE_OPAQUE_WIN32_KMT_KHR); + extMemProperties.push_back((cl_mem_properties)handle); + break; +#else + case VULKAN_EXTERNAL_MEMORY_HANDLE_TYPE_OPAQUE_FD: + fd = (int)vkDeviceMem->getHandle(vkExternalMemoryHandleType); + errNum = check_external_memory_handle_type( + device, CL_EXTERNAL_MEMORY_HANDLE_OPAQUE_FD_KHR); + extMemProperties.push_back( + (cl_mem_properties)CL_EXTERNAL_MEMORY_HANDLE_OPAQUE_FD_KHR); + extMemProperties.push_back((cl_mem_properties)fd); + break; +#endif + default: + errNum = TEST_FAIL; + log_error("Unsupported external memory handle type \n"); + break; + } + if (errNum != CL_SUCCESS) + { + log_error("Checks failed for " + "CL_DEVICE_EXTERNAL_MEMORY_IMPORT_HANDLE_TYPES_KHR\n"); + return TEST_FAIL; + } + extMemProperties.push_back(0); + + const VkImageCreateInfo VulkanImageCreateInfo = + vkImage3D.getVkImageCreateInfo(); + + errNum = getCLImageInfoFromVkImageInfo(&VulkanImageCreateInfo, + totalImageMemSize, &img_format, + &image_desc); + if (errNum != CL_SUCCESS) + { + log_error("getCLImageInfoFromVkImageInfo failed!!!"); + return TEST_FAIL; + } + + clMemWrapper image; + + // Pass valid properties, image_desc and image_format + image = clCreateImageWithProperties( + context, extMemProperties.data(), CL_MEM_READ_WRITE, &img_format, + &image_desc, NULL /* host_ptr */, &errNum); + test_error(errNum, "Unable to create Image with Properties"); + 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, NULL, &errNum); + test_failure_error(errNum, CL_INVALID_IMAGE_FORMAT_DESCRIPTOR, + "Image creation must fail with " + "CL_INVALID_IMAGE_FORMAT_DESCRIPTOR" + "when image desc passed as NULL"); + + image.reset(); + + // Passing image_desc as NULL + image = clCreateImageWithProperties(context, extMemProperties.data(), + CL_MEM_READ_WRITE, &img_format, + NULL, NULL, &errNum); + test_failure_error(errNum, CL_INVALID_IMAGE_DESCRIPTOR, + "Image creation must fail with " + "CL_INVALID_IMAGE_DESCRIPTOR " + "when image desc passed as NULL"); + image.reset(); + + return TEST_PASS; + } +}; + +} // anonymous namespace int test_consistency_external_for_3dimage(cl_device_id deviceID, - cl_context _context, - cl_command_queue _queue, + cl_context context, + cl_command_queue defaultQueue, int num_elements) { - cl_int errNum; - VulkanDevice vkDevice; - - // Context and command queue creation - cl_platform_id platform = NULL; - cl_context context = NULL; - cl_command_queue cmd_queue = NULL; - - cl_context_properties contextProperties[] = { CL_CONTEXT_PLATFORM, 0, 0 }; - errNum = clGetPlatformIDs(1, &platform, NULL); - test_error(errNum, "Failed to get platform id"); - - contextProperties[1] = (cl_context_properties)platform; - - context = clCreateContextFromType(contextProperties, CL_DEVICE_TYPE_GPU, - NULL, NULL, &errNum); - test_error(errNum, "Unable to create context with properties"); - - cmd_queue = clCreateCommandQueue(context, deviceID, 0, &errNum); - test_error(errNum, "Unable to create command queue"); - - cl_device_id devList[] = { deviceID, NULL }; - -#ifdef _WIN32 - if (!is_extension_available(devList[0], "cl_khr_external_memory_win32")) - { - throw std::runtime_error("Device does not support" - "cl_khr_external_memory_win32 extension \n"); - } -#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"); - } -#endif - uint32_t width = 256; - uint32_t height = 16; - uint32_t depth = 10; - cl_image_desc image_desc; - memset(&image_desc, 0x0, sizeof(cl_image_desc)); - cl_image_format img_format = { 0 }; - - VulkanExternalMemoryHandleType vkExternalMemoryHandleType = - getSupportedVulkanExternalMemoryHandleTypeList()[0]; - - VulkanImageTiling vulkanImageTiling = - vkClExternalMemoryHandleTilingAssumption( - deviceID, vkExternalMemoryHandleType, &errNum); - ASSERT_SUCCESS(errNum, "Failed to query OpenCL tiling mode"); - - VulkanImage3D vkImage3D = - VulkanImage3D(vkDevice, VULKAN_FORMAT_R8G8B8A8_UNORM, width, height, - depth, vulkanImageTiling, 1, vkExternalMemoryHandleType); - - const VulkanMemoryTypeList& memoryTypeList = vkImage3D.getMemoryTypeList(); - uint64_t totalImageMemSize = vkImage3D.getSize(); - - log_info("Memory type index: %u\n", (uint32_t)memoryTypeList[0]); - log_info("Memory type property: %d\n", - memoryTypeList[0].getMemoryTypeProperty()); - log_info("Image size : %lu\n", totalImageMemSize); - - VulkanDeviceMemory* vkDeviceMem = new VulkanDeviceMemory( - vkDevice, vkImage3D, memoryTypeList[0], vkExternalMemoryHandleType); - vkDeviceMem->bindImage(vkImage3D, 0); - - void* handle = NULL; - int fd; - std::vector extMemProperties{ - (cl_mem_properties)CL_MEM_DEVICE_HANDLE_LIST_KHR, - (cl_mem_properties)devList[0], - (cl_mem_properties)CL_MEM_DEVICE_HANDLE_LIST_END_KHR, - }; - switch (vkExternalMemoryHandleType) - { -#ifdef _WIN32 - case VULKAN_EXTERNAL_MEMORY_HANDLE_TYPE_OPAQUE_WIN32_NT: - handle = vkDeviceMem->getHandle(vkExternalMemoryHandleType); - errNum = check_external_memory_handle_type( - devList[0], CL_EXTERNAL_MEMORY_HANDLE_OPAQUE_WIN32_KHR); - extMemProperties.push_back( - (cl_mem_properties)CL_EXTERNAL_MEMORY_HANDLE_OPAQUE_WIN32_KHR); - extMemProperties.push_back((cl_mem_properties)handle); - break; - case VULKAN_EXTERNAL_MEMORY_HANDLE_TYPE_OPAQUE_WIN32_KMT: - handle = vkDeviceMem->getHandle(vkExternalMemoryHandleType); - errNum = check_external_memory_handle_type( - devList[0], CL_EXTERNAL_MEMORY_HANDLE_OPAQUE_WIN32_KMT_KHR); - extMemProperties.push_back( - (cl_mem_properties) - CL_EXTERNAL_MEMORY_HANDLE_OPAQUE_WIN32_KMT_KHR); - extMemProperties.push_back((cl_mem_properties)handle); - break; -#else - case VULKAN_EXTERNAL_MEMORY_HANDLE_TYPE_OPAQUE_FD: - fd = (int)vkDeviceMem->getHandle(vkExternalMemoryHandleType); - errNum = check_external_memory_handle_type( - devList[0], CL_EXTERNAL_MEMORY_HANDLE_OPAQUE_FD_KHR); - extMemProperties.push_back( - (cl_mem_properties)CL_EXTERNAL_MEMORY_HANDLE_OPAQUE_FD_KHR); - extMemProperties.push_back((cl_mem_properties)fd); - break; -#endif - default: - errNum = TEST_FAIL; - log_error("Unsupported external memory handle type \n"); - break; - } - if (errNum != CL_SUCCESS) - { - log_error("Checks failed for " - "CL_DEVICE_EXTERNAL_MEMORY_IMPORT_HANDLE_TYPES_KHR\n"); - return TEST_FAIL; - } - extMemProperties.push_back(0); - - const VkImageCreateInfo VulkanImageCreateInfo = - vkImage3D.getVkImageCreateInfo(); - - errNum = getCLImageInfoFromVkImageInfo( - &VulkanImageCreateInfo, totalImageMemSize, &img_format, &image_desc); - if (errNum != CL_SUCCESS) - { - log_error("getCLImageInfoFromVkImageInfo failed!!!"); - return TEST_FAIL; - } - - clMemWrapper image; - - // Pass valid properties, image_desc and image_format - image = clCreateImageWithProperties( - context, extMemProperties.data(), CL_MEM_READ_WRITE, &img_format, - &image_desc, NULL /* host_ptr */, &errNum); - test_error(errNum, "Unable to create Image with Properties"); - 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, - NULL, &errNum); - test_failure_error(errNum, CL_INVALID_IMAGE_FORMAT_DESCRIPTOR, - "Image creation must fail with " - "CL_INVALID_IMAGE_FORMAT_DESCRIPTOR" - "when image desc passed as NULL"); - - image.reset(); - - // Passing image_desc as NULL - image = clCreateImageWithProperties(context, extMemProperties.data(), - CL_MEM_READ_WRITE, &img_format, NULL, - NULL, &errNum); - test_failure_error(errNum, CL_INVALID_IMAGE_DESCRIPTOR, - "Image creation must fail with " - "CL_INVALID_IMAGE_DESCRIPTOR " - "when image desc passed as NULL"); - image.reset(); - - if (cmd_queue) clReleaseCommandQueue(cmd_queue); - if (context) clReleaseContext(context); - - return TEST_PASS; + return MakeAndRunTest( + deviceID, context, defaultQueue, num_elements); } diff --git a/test_conformance/vulkan/test_vulkan_interop_buffer.cpp b/test_conformance/vulkan/test_vulkan_interop_buffer.cpp index 56fd485c..e5f1a728 100644 --- a/test_conformance/vulkan/test_vulkan_interop_buffer.cpp +++ b/test_conformance/vulkan/test_vulkan_interop_buffer.cpp @@ -1,5 +1,5 @@ // -// Copyright (c) 2022 The Khronos Group Inc. +// Copyright (c) 2024 The Khronos Group Inc. // // Licensed under the Apache License, Version 2.0 (the "License"); // you may not use this file except in compliance with the License. @@ -24,62 +24,66 @@ #include #include "harness/errorHelpers.h" #include "harness/os_helpers.h" -#include "deviceInfo.h" + +#include "vulkan_test_base.h" +#include "opencl_vulkan_wrapper.hpp" #define MAX_BUFFERS 5 #define MAX_IMPORTS 5 #define BUFFERSIZE 3000 -static cl_uchar uuid[CL_UUID_SIZE_KHR]; -static cl_device_id deviceId = NULL; namespace { + +cl_uchar uuid[CL_UUID_SIZE_KHR]; +cl_device_id deviceId = nullptr; + struct Params { uint32_t numBuffers; uint32_t bufferSize; uint32_t interBufferOffset; }; -} const char *kernel_text_numbuffer_1 = " \ -__kernel void clUpdateBuffer(int bufferSize, __global unsigned char *a) { \n\ - int gid = get_global_id(0); \n\ - if (gid < bufferSize) { \n\ - a[gid]++; \n\ - } \n\ -}"; + __kernel void clUpdateBuffer(int bufferSize, __global unsigned char *a) { \n\ + int gid = get_global_id(0); \n\ + if (gid < bufferSize) { \n\ + a[gid]++; \n\ + } \n\ + }"; const char *kernel_text_numbuffer_2 = " \ -__kernel void clUpdateBuffer(int bufferSize, __global unsigned char *a, __global unsigned char *b) { \n\ - int gid = get_global_id(0); \n\ - if (gid < bufferSize) { \n\ - a[gid]++; \n\ - b[gid]++;\n\ - } \n\ -}"; + __kernel void clUpdateBuffer(int bufferSize, __global unsigned char *a, __global unsigned char *b) { \n\ + int gid = get_global_id(0); \n\ + if (gid < bufferSize) { \n\ + a[gid]++; \n\ + b[gid]++;\n\ + } \n\ + }"; const char *kernel_text_numbuffer_4 = " \ -__kernel void clUpdateBuffer(int bufferSize, __global unsigned char *a, __global unsigned char *b, __global unsigned char *c, __global unsigned char *d) { \n\ - int gid = get_global_id(0); \n\ - if (gid < bufferSize) { \n\ - a[gid]++;\n\ - b[gid]++; \n\ - c[gid]++; \n\ - d[gid]++; \n\ - } \n\ -}"; + __kernel void clUpdateBuffer(int bufferSize, __global unsigned char *a, __global unsigned char *b, __global unsigned char *c, __global unsigned char *d) { \n\ + int gid = get_global_id(0); \n\ + if (gid < bufferSize) { \n\ + a[gid]++;\n\ + b[gid]++; \n\ + c[gid]++; \n\ + d[gid]++; \n\ + } \n\ + }"; const char *kernel_text_verify = " \ -__kernel void checkKernel(__global unsigned char *ptr, int size, int expVal, __global unsigned char *err) \n\ -{ \n\ - int idx = get_global_id(0); \n\ - if ((idx < size) && (*err == 0)) { \n\ - if (ptr[idx] != expVal){ \n\ - *err = 1; \n\ - } \n\ - } \n\ -}"; + __kernel void checkKernel(__global unsigned char *ptr, int size, int expVal, __global unsigned char *err) \n\ + { \n\ + int idx = get_global_id(0); \n\ + if ((idx < size) && (*err == 0)) { \n\ + if (ptr[idx] != expVal){ \n\ + *err = 1; \n\ + } \n\ + } \n\ + }"; + int run_test_with_two_queue( cl_context &context, cl_command_queue &cmd_queue1, @@ -114,7 +118,8 @@ int run_test_with_two_queue( VulkanSemaphore vkCl2VkSemaphore(vkDevice, vkExternalSemaphoreHandleType); std::shared_ptr fence = nullptr; - VulkanQueue &vkQueue = vkDevice.getQueue(); + VulkanQueue &vkQueue = + vkDevice.getQueue(getVulkanQueueFamily(vkDevice.getPhysicalDevice())); std::vector vkBufferShader = readFile("buffer.spv", exe_dir()); @@ -150,6 +155,7 @@ int run_test_with_two_queue( } const uint32_t maxIter = innerIterations; + VulkanCommandPool vkCommandPool(vkDevice); VulkanCommandBuffer vkCommandBuffer(vkDevice, vkCommandPool); @@ -446,7 +452,8 @@ int run_test_with_one_queue( VulkanSemaphore vkCl2VkSemaphore(vkDevice, vkExternalSemaphoreHandleType); std::shared_ptr fence = nullptr; - VulkanQueue &vkQueue = vkDevice.getQueue(); + VulkanQueue &vkQueue = + vkDevice.getQueue(getVulkanQueueFamily(vkDevice.getPhysicalDevice())); std::vector vkBufferShader = readFile("buffer.spv", exe_dir()); @@ -482,6 +489,7 @@ int run_test_with_one_queue( } const uint32_t maxIter = innerIterations; + VulkanCommandPool vkCommandPool(vkDevice); VulkanCommandBuffer vkCommandBuffer(vkDevice, vkCommandPool); @@ -749,7 +757,7 @@ int run_test_with_multi_import_same_ctx( VulkanSemaphore vkCl2VkSemaphore(vkDevice, vkExternalSemaphoreHandleType); std::shared_ptr fence = nullptr; - VulkanQueue &vkQueue = vkDevice.getQueue(); + VulkanQueue &vkQueue = vkDevice.getQueue(getVulkanQueueFamily()); std::vector vkBufferShader = readFile("buffer.spv", exe_dir()); @@ -820,221 +828,213 @@ int run_test_with_multi_import_same_ctx( memoryType.getMemoryTypeProperty()); - cl_mem buffers[MAX_BUFFERS][MAX_IMPORTS]; - VulkanBufferList vkBufferList(numBuffers, vkDevice, bufferSize, - vkExternalMemoryHandleType); + cl_mem buffers[MAX_BUFFERS][MAX_IMPORTS]; + VulkanBufferList vkBufferList(numBuffers, vkDevice, bufferSize, + vkExternalMemoryHandleType); - for (size_t bIdx = 0; bIdx < numBuffers; bIdx++) + for (size_t bIdx = 0; bIdx < numBuffers; bIdx++) + { + vkBufferListDeviceMemory.push_back(new VulkanDeviceMemory( + vkDevice, vkBufferList[bIdx], memoryType, + vkExternalMemoryHandleType)); + + std::vector pExternalMemory; + for (size_t cl_bIdx = 0; cl_bIdx < numImports; cl_bIdx++) { - 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, bufferSize, context, - deviceId)); - } - externalMemory.push_back(pExternalMemory); + pExternalMemory.push_back( + new clExternalMemory(vkBufferListDeviceMemory[bIdx], + vkExternalMemoryHandleType, + bufferSize, context, deviceId)); } + externalMemory.push_back(pExternalMemory); + } - clFinish(cmd_queue1); - Params *params = (Params *)vkParamsDeviceMemory.map(); - params->numBuffers = numBuffers; - params->bufferSize = bufferSize; - params->interBufferOffset = 0; - vkParamsDeviceMemory.unmap(); - vkDescriptorSet.update(0, vkParamsBuffer); - for (size_t bIdx = 0; bIdx < vkBufferList.size(); bIdx++) + clFinish(cmd_queue1); + Params *params = (Params *)vkParamsDeviceMemory.map(); + params->numBuffers = numBuffers; + 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], + 0); + for (size_t cl_bIdx = 0; cl_bIdx < numImports; cl_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++) - { - buffers[bIdx][cl_bIdx] = - externalMemory[bIdx][cl_bIdx] - ->getExternalMemoryBuffer(); - } + buffers[bIdx][cl_bIdx] = externalMemory[bIdx][cl_bIdx] + ->getExternalMemoryBuffer(); } - vkDescriptorSet.updateArray(1, numBuffers, vkBufferList); - vkCommandBuffer.begin(); - vkCommandBuffer.bindPipeline(vkComputePipeline); - vkCommandBuffer.bindDescriptorSets( - vkComputePipeline, vkPipelineLayout, vkDescriptorSet); - vkCommandBuffer.dispatch(512, 1, 1); - vkCommandBuffer.end(); + } + vkDescriptorSet.updateArray(1, numBuffers, vkBufferList); + vkCommandBuffer.begin(); + vkCommandBuffer.bindPipeline(vkComputePipeline); + vkCommandBuffer.bindDescriptorSets( + vkComputePipeline, vkPipelineLayout, vkDescriptorSet); + vkCommandBuffer.dispatch(512, 1, 1); + vkCommandBuffer.end(); - 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] = bufferSize; + 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] = bufferSize; - for (uint32_t iter = 0; iter < maxIter; iter++) + for (uint32_t iter = 0; iter < maxIter; iter++) + { + if (use_fence) { - 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); - } - } - - 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"); - } - - for (uint8_t launchIter = 0; launchIter < numImports; - launchIter++) - { - err = clSetKernelArg(update_buffer_kernel, 0, - sizeof(uint32_t), - (void *)&bufferSize); - for (int i = 0; i < numBuffers; i++) - { - err |= clSetKernelArg( - update_buffer_kernel, i + 1, sizeof(cl_mem), - (void *)&(buffers[i][launchIter])); - err = clEnqueueAcquireExternalMemObjectsKHRptr( - cmd_queue1, 1, &buffers[i][launchIter], 0, - nullptr, nullptr); - test_error_and_cleanup(err, CLEANUP, - "Failed to acquire buffers"); - } - test_error_and_cleanup( - err, CLEANUP, - "Error: Failed to set arg values for " - "kernel\n "); - - err = clEnqueueNDRangeKernel( - 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 "); - - for (int i = 0; i < numBuffers; i++) - { - err = clEnqueueReleaseExternalMemObjectsKHRptr( - cmd_queue1, 1, &buffers[i][launchIter], 0, - nullptr, nullptr); - test_error_and_cleanup(err, CLEANUP, - "Failed to release buffers"); - } - } - 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"); + vkQueue.submit(vkCl2VkSemaphore, vkCommandBuffer, + vkVk2CLSemaphore); } } - error_2 = (uint8_t *)malloc(sizeof(uint8_t)); - if (NULL == error_2) + if (use_fence) { - test_fail_and_cleanup(err, CLEANUP, - "Not able to allocate memory\n"); + fence->wait(); + } + else + { + err = clVk2CLExternalSemaphore->wait(cmd_queue1); + test_error_and_cleanup(err, CLEANUP, + "Error: failed to wait on " + "CL external semaphore\n"); } - error_1 = clCreateBuffer(context, CL_MEM_WRITE_ONLY, - sizeof(uint8_t), NULL, &err); - 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); - test_error_and_cleanup(err, CLEANUP, - "Error: clEnqueueWriteBuffer \n"); - - calc_max_iter = maxIter * (numImports + 1); - - for (int i = 0; i < vkBufferList.size(); i++) + for (uint8_t launchIter = 0; launchIter < numImports; + launchIter++) { - err = clSetKernelArg(verify_kernel, 0, sizeof(cl_mem), - (void *)&(buffers[i][0])); - err |= clSetKernelArg(verify_kernel, 1, sizeof(int), - &bufferSize); - err |= clSetKernelArg(verify_kernel, 2, sizeof(int), - &calc_max_iter); - err |= clSetKernelArg(verify_kernel, 3, sizeof(cl_mem), - (void *)&error_1); + err = clSetKernelArg(update_buffer_kernel, 0, + sizeof(uint32_t), (void *)&bufferSize); + for (int i = 0; i < numBuffers; i++) + { + err |= clSetKernelArg( + update_buffer_kernel, i + 1, sizeof(cl_mem), + (void *)&(buffers[i][launchIter])); + err = clEnqueueAcquireExternalMemObjectsKHRptr( + cmd_queue1, 1, &buffers[i][launchIter], 0, nullptr, + nullptr); + test_error_and_cleanup(err, CLEANUP, + "Failed to acquire buffers"); + } test_error_and_cleanup( err, CLEANUP, "Error: Failed to set arg values for " - "verify_kernel \n"); + "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"); + err = clEnqueueNDRangeKernel( + 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 "); - err = clEnqueueReadBuffer(cmd_queue1, error_1, CL_TRUE, 0, - sizeof(uint8_t), error_2, 0, NULL, - NULL); - test_error_and_cleanup( - err, CLEANUP, "Error: Failed read output, error \n"); - - if (*error_2 == 1) + for (int i = 0; i < numBuffers; i++) { - test_fail_and_cleanup( - err, CLEANUP, - " vulkan_opencl_buffer test FAILED\n"); + err = clEnqueueReleaseExternalMemObjectsKHRptr( + cmd_queue1, 1, &buffers[i][launchIter], 0, nullptr, + nullptr); + test_error_and_cleanup(err, CLEANUP, + "Failed to release buffers"); } } - for (size_t i = 0; i < vkBufferList.size(); i++) + if (use_fence) { - for (size_t j = 0; j < numImports; j++) - { - delete externalMemory[i][j]; - } + clFinish(cmd_queue1); } - for (size_t i = 0; i < vkBufferListDeviceMemory.size(); i++) + else if (!use_fence && iter != (maxIter - 1)) { - delete vkBufferListDeviceMemory[i]; + err = clCl2VkExternalSemaphore->signal(cmd_queue1); + test_error_and_cleanup(err, CLEANUP, + "Failed to signal CL semaphore\n"); } - vkBufferListDeviceMemory.erase(vkBufferListDeviceMemory.begin(), - vkBufferListDeviceMemory.end()); - for (size_t i = 0; i < externalMemory.size(); i++) + } + + error_2 = (uint8_t *)malloc(sizeof(uint8_t)); + if (NULL == error_2) + { + 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); + 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); + 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), &bufferSize); + err |= clSetKernelArg(verify_kernel, 2, sizeof(int), + &calc_max_iter); + err |= clSetKernelArg(verify_kernel, 3, sizeof(cl_mem), + (void *)&error_1); + 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"); + + err = clEnqueueReadBuffer(cmd_queue1, error_1, CL_TRUE, 0, + sizeof(uint8_t), error_2, 0, NULL, + NULL); + test_error_and_cleanup(err, CLEANUP, + "Error: Failed read output, error \n"); + + if (*error_2 == 1) { - externalMemory[i].erase(externalMemory[i].begin(), - externalMemory[i].begin() - + numBuffers); + test_fail_and_cleanup( + err, CLEANUP, " vulkan_opencl_buffer test FAILED\n"); } - externalMemory.clear(); + } + for (size_t i = 0; i < vkBufferList.size(); i++) + { + for (size_t j = 0; j < numImports; j++) + { + delete externalMemory[i][j]; + } + } + for (size_t i = 0; i < vkBufferListDeviceMemory.size(); i++) + { + delete vkBufferListDeviceMemory[i]; + } + vkBufferListDeviceMemory.erase(vkBufferListDeviceMemory.begin(), + vkBufferListDeviceMemory.end()); + for (size_t i = 0; i < externalMemory.size(); i++) + { + externalMemory[i].erase(externalMemory[i].begin(), + externalMemory[i].begin() + numBuffers); + } + externalMemory.clear(); } } CLEANUP: @@ -1097,7 +1097,7 @@ int run_test_with_multi_import_diff_ctx( VulkanSemaphore vkCl2VkSemaphore(vkDevice, vkExternalSemaphoreHandleType); std::shared_ptr fence = nullptr; - VulkanQueue &vkQueue = vkDevice.getQueue(); + VulkanQueue &vkQueue = vkDevice.getQueue(getVulkanQueueFamily()); std::vector vkBufferShader = readFile("buffer.spv", exe_dir()); @@ -1273,9 +1273,9 @@ int run_test_with_multi_import_diff_ctx( else { err = clVk2CLExternalSemaphore->wait(cmd_queue1); - test_error_and_cleanup( - err, CLEANUP, - "Error: failed to wait on CL external semaphore\n"); + test_error_and_cleanup(err, CLEANUP, + "Error: failed to wait on " + "CL external semaphore\n"); } for (uint8_t launchIter = 0; launchIter < numImports; @@ -1354,204 +1354,193 @@ int run_test_with_multi_import_diff_ctx( } } - if (use_fence) - { - fence->wait(); - } - else - { - err = clVk2CLExternalSemaphore2->wait(cmd_queue2); - test_error_and_cleanup( - err, CLEANUP, - "Error: failed to wait on CL external semaphore\n"); - } + if (use_fence) + { + fence->wait(); + } + else + { + 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; - launchIter++) + for (uint8_t launchIter = 0; launchIter < numImports; + launchIter++) + { + err = clSetKernelArg(update_buffer_kernel2[launchIter], 0, + sizeof(uint32_t), (void *)&bufferSize); + test_error_and_cleanup(err, CLEANUP, + "Failed to set kernel arg"); + + for (int i = 0; i < numBuffers; i++) { - err = clSetKernelArg(update_buffer_kernel2[launchIter], - 0, sizeof(uint32_t), - (void *)&bufferSize); + err = clSetKernelArg( + update_buffer_kernel2[launchIter], i + 1, + sizeof(cl_mem), (void *)&(buffers2[i][launchIter])); test_error_and_cleanup(err, CLEANUP, "Failed to set kernel arg"); - for (int i = 0; i < numBuffers; i++) - { - err = clSetKernelArg( - update_buffer_kernel2[launchIter], i + 1, - sizeof(cl_mem), - (void *)&(buffers2[i][launchIter])); - test_error_and_cleanup(err, CLEANUP, - "Failed to set kernel arg"); + err = clEnqueueAcquireExternalMemObjectsKHRptr( + cmd_queue2, 1, &buffers2[i][launchIter], 0, nullptr, + nullptr); + test_error_and_cleanup(err, CLEANUP, + "Failed to acquire buffers"); + } + test_error_and_cleanup( + err, CLEANUP, + "Error: Failed to set arg values for " + "kernel\n "); - err = clEnqueueAcquireExternalMemObjectsKHRptr( - cmd_queue2, 1, &buffers2[i][launchIter], 0, - nullptr, nullptr); - test_error_and_cleanup(err, CLEANUP, - "Failed to acquire buffers"); - } - test_error_and_cleanup( - err, CLEANUP, - "Error: Failed to set arg values for " - "kernel\n "); - - err = clEnqueueNDRangeKernel( - cmd_queue2, update_buffer_kernel2[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 "); - for (int i = 0; i < numBuffers; i++) - { - err = clEnqueueReleaseExternalMemObjectsKHRptr( - cmd_queue2, 1, &buffers2[i][launchIter], 0, - nullptr, nullptr); - test_error_and_cleanup(err, CLEANUP, - "Failed to release buffers"); - } - } - if (use_fence) + err = clEnqueueNDRangeKernel( + cmd_queue2, update_buffer_kernel2[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 "); + for (int i = 0; i < numBuffers; i++) { - clFinish(cmd_queue2); - } - else if (!use_fence && iter != (maxIter - 1)) - { - err = clCl2VkExternalSemaphore2->signal(cmd_queue2); - test_error_and_cleanup( - err, CLEANUP, "Failed to signal CL semaphore\n"); + err = clEnqueueReleaseExternalMemObjectsKHRptr( + cmd_queue2, 1, &buffers2[i][launchIter], 0, nullptr, + nullptr); + test_error_and_cleanup(err, CLEANUP, + "Failed to release buffers"); } + } + if (use_fence) + { + clFinish(cmd_queue2); + } + else if (!use_fence && iter != (maxIter - 1)) + { + 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) + { + test_fail_and_cleanup(err, CLEANUP, + "Not able to allocate memory\n"); } - clFinish(cmd_queue2); - error_3 = (uint8_t *)malloc(sizeof(uint8_t)); - if (NULL == error_3) - { - 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); + error_1 = clCreateBuffer(context, CL_MEM_WRITE_ONLY, + sizeof(uint8_t), NULL, &err); + test_error_and_cleanup(err, CLEANUP, "Error: clCreateBuffer \n"); + + error_2 = clCreateBuffer(context2, CL_MEM_WRITE_ONLY, + sizeof(uint8_t), NULL, &err); + 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); + 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); + 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++) + { + err = clSetKernelArg(verify_kernel, 0, sizeof(cl_mem), + (void *)&(buffers1[i][0])); + err |= + clSetKernelArg(verify_kernel, 1, sizeof(int), &pBufferSize); + err |= clSetKernelArg(verify_kernel, 2, sizeof(int), + &calc_max_iter); + err |= clSetKernelArg(verify_kernel, 3, sizeof(cl_mem), + (void *)&error_1); test_error_and_cleanup(err, CLEANUP, - "Error: clCreateBuffer \n"); + "Error: Failed to set arg values for " + "verify_kernel \n"); - error_2 = clCreateBuffer(context2, CL_MEM_WRITE_ONLY, - sizeof(uint8_t), NULL, &err); + err = clEnqueueNDRangeKernel(cmd_queue1, verify_kernel, 1, NULL, + global_work_size, NULL, 0, NULL, + NULL); test_error_and_cleanup(err, CLEANUP, - "Error: clCreateBuffer \n"); + "Error: Failed to launch verify_kernel," + "error\n"); - uint8_t val = 0; - err = - clEnqueueWriteBuffer(cmd_queue1, error_1, CL_TRUE, 0, - sizeof(uint8_t), &val, 0, NULL, NULL); + err = clEnqueueReadBuffer(cmd_queue1, error_1, CL_TRUE, 0, + sizeof(uint8_t), error_3, 0, NULL, + NULL); test_error_and_cleanup(err, CLEANUP, - "Error: Failed read output, error \n"); + "Error: Failed read output, error\n"); - err = - clEnqueueWriteBuffer(cmd_queue2, error_2, CL_TRUE, 0, - sizeof(uint8_t), &val, 0, NULL, NULL); + if (*error_3 == 1) + { + test_fail_and_cleanup( + err, CLEANUP, + "&&&& vulkan_opencl_buffer test FAILED\n"); + } + } + *error_3 = 0; + for (int i = 0; i < vkBufferList.size(); i++) + { + err = clSetKernelArg(verify_kernel2, 0, sizeof(cl_mem), + (void *)&(buffers2[i][0])); + err |= clSetKernelArg(verify_kernel2, 1, sizeof(int), + &pBufferSize); + err |= clSetKernelArg(verify_kernel2, 2, sizeof(int), + &calc_max_iter); + err |= clSetKernelArg(verify_kernel2, 3, sizeof(cl_mem), + (void *)&error_2); test_error_and_cleanup(err, CLEANUP, - "Error: Failed read output, error \n"); + "Error: Failed to set arg values for " + "verify_kernel \n"); - calc_max_iter = maxIter * 2 * (numBuffers + 1); - for (int i = 0; i < numBuffers; i++) + err = clEnqueueNDRangeKernel(cmd_queue2, verify_kernel2, 1, + NULL, global_work_size, NULL, 0, + NULL, NULL); + 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); + test_error_and_cleanup(err, CLEANUP, + "Error: Failed read output, error\n"); + + if (*error_3 == 1) { - err = clSetKernelArg(verify_kernel, 0, sizeof(cl_mem), - (void *)&(buffers1[i][0])); - err |= clSetKernelArg(verify_kernel, 1, sizeof(int), - &pBufferSize); - err |= clSetKernelArg(verify_kernel, 2, sizeof(int), - &calc_max_iter); - err |= clSetKernelArg(verify_kernel, 3, sizeof(cl_mem), - (void *)&error_1); - test_error_and_cleanup( + test_fail_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"); - - err = clEnqueueReadBuffer(cmd_queue1, error_1, CL_TRUE, 0, - sizeof(uint8_t), error_3, 0, NULL, - NULL); - test_error_and_cleanup( - err, CLEANUP, "Error: Failed read output, error\n"); - - if (*error_3 == 1) - { - test_fail_and_cleanup( - err, CLEANUP, - "&&&& vulkan_opencl_buffer test FAILED\n"); - } + "&&&& vulkan_opencl_buffer test FAILED\n"); } - *error_3 = 0; - for (int i = 0; i < vkBufferList.size(); i++) + } + for (size_t i = 0; i < vkBufferList.size(); i++) + { + for (size_t j = 0; j < numImports; j++) { - err = clSetKernelArg(verify_kernel2, 0, sizeof(cl_mem), - (void *)&(buffers2[i][0])); - err |= clSetKernelArg(verify_kernel2, 1, sizeof(int), - &pBufferSize); - err |= clSetKernelArg(verify_kernel2, 2, sizeof(int), - &calc_max_iter); - err |= clSetKernelArg(verify_kernel2, 3, sizeof(cl_mem), - (void *)&error_2); - 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); - 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); - test_error_and_cleanup( - err, CLEANUP, "Error: Failed read output, error\n"); - - if (*error_3 == 1) - { - test_fail_and_cleanup( - err, CLEANUP, - "&&&& vulkan_opencl_buffer test FAILED\n"); - } + delete externalMemory1[i][j]; + delete externalMemory2[i][j]; } - for (size_t i = 0; i < vkBufferList.size(); i++) - { - for (size_t j = 0; j < numImports; j++) - { - delete externalMemory1[i][j]; - delete externalMemory2[i][j]; - } - } - for (size_t i = 0; i < vkBufferListDeviceMemory.size(); i++) - { - delete vkBufferListDeviceMemory[i]; - } - vkBufferListDeviceMemory.erase(vkBufferListDeviceMemory.begin(), - vkBufferListDeviceMemory.end()); - for (size_t i = 0; i < externalMemory1.size(); i++) - { - externalMemory1[i].erase(externalMemory1[i].begin(), - externalMemory1[i].begin() - + numBuffers); - externalMemory2[i].erase(externalMemory2[i].begin(), - externalMemory2[i].begin() - + numBuffers); - } - externalMemory1.clear(); - externalMemory2.clear(); + } + for (size_t i = 0; i < vkBufferListDeviceMemory.size(); i++) + { + delete vkBufferListDeviceMemory[i]; + } + vkBufferListDeviceMemory.erase(vkBufferListDeviceMemory.begin(), + vkBufferListDeviceMemory.end()); + for (size_t i = 0; i < externalMemory1.size(); i++) + { + externalMemory1[i].erase(externalMemory1[i].begin(), + externalMemory1[i].begin() + + numBuffers); + externalMemory2[i].erase(externalMemory2[i].begin(), + externalMemory2[i].begin() + + numBuffers); + } + externalMemory1.clear(); + externalMemory2.clear(); } } CLEANUP: @@ -1597,80 +1586,52 @@ CLEANUP: return err; } -int test_buffer_common(cl_device_id device_, cl_context context_, - cl_command_queue queue_, int numElements_, - bool use_fence) + +struct BufferTestBase : public VulkanTestBase { + BufferTestBase(cl_device_id device, cl_context context, + cl_command_queue queue, cl_int nelems) + : VulkanTestBase(device, context, queue, nelems) + {} - int current_device = 0; - int device_count = 0; - int devices_prohibited = 0; - cl_int errNum = CL_SUCCESS; - cl_platform_id platform = NULL; - size_t extensionSize = 0; - cl_uint num_devices = 0; - cl_uint device_no = 0; - const size_t bufsize = BUFFERSIZE; - char buf[BUFFERSIZE]; - cl_device_id *devices; - char *extensions = NULL; - cl_kernel verify_kernel; - cl_kernel verify_kernel2; - cl_kernel kernel[3] = { NULL, NULL, NULL }; - cl_kernel kernel2[3] = { NULL, NULL, NULL }; - const char *program_source_const[3] = { kernel_text_numbuffer_1, - kernel_text_numbuffer_2, - kernel_text_numbuffer_4 }; - const char *program_source_const_verify; - size_t program_source_length; - cl_command_queue cmd_queue1 = NULL; - cl_command_queue cmd_queue2 = NULL; - cl_command_queue cmd_queue3 = NULL; - cl_context context = NULL; - cl_program program[3] = { NULL, NULL, NULL }; - cl_program program_verify, program_verify2; - cl_context context2 = NULL; - - - VulkanDevice vkDevice; - uint32_t numBuffersList[] = { 1, 2, 4 }; - uint32_t bufferSizeList[] = { 4 * 1024, 64 * 1024, 2 * 1024 * 1024 }; - uint32_t bufferSizeListforOffset[] = { 256, 512, 1024 }; - - cl_context_properties contextProperties[] = { CL_CONTEXT_PLATFORM, 0, 0 }; - std::vector supportedSemaphoreTypes; - - errNum = clGetPlatformIDs(1, &platform, NULL); - test_error_and_cleanup(errNum, CLEANUP, "Error: Failed to get platform\n"); - - errNum = - clGetDeviceIDs(platform, CL_DEVICE_TYPE_GPU, 0, NULL, &num_devices); - 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) + int test_buffer_common(bool use_fence) { - test_fail_and_cleanup(errNum, CLEANUP, - "Unable to allocate memory for devices\n"); - } - errNum = clGetDeviceIDs(platform, CL_DEVICE_TYPE_GPU, num_devices, devices, - NULL); - test_error_and_cleanup(errNum, CLEANUP, "Failed to get deviceID.\n"); + int current_device = 0; + int device_count = 0; + int devices_prohibited = 0; + cl_int errNum = CL_SUCCESS; + size_t extensionSize = 0; + const size_t bufsize = BUFFERSIZE; + char buf[BUFFERSIZE]; + char *extensions = NULL; + clKernelWrapper verify_kernel; + clKernelWrapper verify_kernel2; + clKernelWrapper kernel[3] = { NULL, NULL, NULL }; + clKernelWrapper kernel2[3] = { NULL, NULL, NULL }; + const char *program_source_const[3] = { kernel_text_numbuffer_1, + kernel_text_numbuffer_2, + kernel_text_numbuffer_4 }; + const char *program_source_const_verify; + size_t program_source_length; + clCommandQueueWrapper cmd_queue1; + clCommandQueueWrapper cmd_queue2; + clCommandQueueWrapper cmd_queue3; - 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_UUID_KHR, - CL_UUID_SIZE_KHR, uuid, NULL); - test_error_and_cleanup(errNum, CLEANUP, "clGetDeviceInfo failed\n"); + clProgramWrapper program[3] = { NULL, NULL, NULL }; + clProgramWrapper program_verify, program_verify2; + clContextWrapper context2; + + uint32_t numBuffersList[] = { 1, 2, 4 }; + uint32_t bufferSizeList[] = { 4 * 1024, 64 * 1024, 2 * 1024 * 1024 }; + uint32_t bufferSizeListforOffset[] = { 256, 512, 1024 }; + + std::vector supportedSemaphoreTypes; if (!use_fence) { supportedSemaphoreTypes = - getSupportedInteropExternalSemaphoreHandleTypes( - devices[device_no], vkDevice); + getSupportedInteropExternalSemaphoreHandleTypes(device, + *vkDevice); } else { @@ -1678,176 +1639,247 @@ int test_buffer_common(cl_device_id device_, cl_context context_, VULKAN_EXTERNAL_SEMAPHORE_HANDLE_TYPE_NONE); } - // If device does not support any semaphores, try the next one if (!use_fence && supportedSemaphoreTypes.empty()) { - continue; + return TEST_FAIL; } - errNum = - memcmp(uuid, vkDevice.getPhysicalDevice().getUUID(), VK_UUID_SIZE); - if (errNum == 0) + if (!use_fence && supportedSemaphoreTypes.empty()) { - break; + test_error_fail( + errNum, "No devices found that support OpenCL semaphores\n"); } - } - if (!use_fence && supportedSemaphoreTypes.empty()) - { - test_fail_and_cleanup( - errNum, CLEANUP, - "No devices found that support OpenCL semaphores\n"); - } + deviceId = device; + cmd_queue1 = clCreateCommandQueue(context, device, 0, &errNum); + test_error(errNum, "Error: Failed to create command queue!\n"); - if (device_no >= num_devices) - { - 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); - 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); - test_error_and_cleanup(errNum, CLEANUP, - "Error: Failed to create command queue!\n"); - - cmd_queue2 = clCreateCommandQueue(context, devices[device_no], 0, &errNum); - 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++) - { - program_source_length = strlen(program_source_const[i]); - program[i] = - clCreateProgramWithSource(context, 1, &program_source_const[i], - &program_source_length, &errNum); - errNum = clBuildProgram(program[i], 0, NULL, NULL, NULL, NULL); - test_error_and_cleanup(errNum, CLEANUP, - "Error: Failed to build program \n"); - - // create the kernel - kernel[i] = clCreateKernel(program[i], "clUpdateBuffer", &errNum); - test_error_and_cleanup(errNum, CLEANUP, "clCreateKernel failed \n"); - } - - program_source_const_verify = kernel_text_verify; - program_source_length = strlen(program_source_const_verify); - program_verify = - clCreateProgramWithSource(context, 1, &program_source_const_verify, - &program_source_length, &errNum); - errNum = clBuildProgram(program_verify, 0, NULL, NULL, NULL, NULL); - test_error_and_cleanup(errNum, CLEANUP, - "Error: Failed to build program2\n"); - - verify_kernel = clCreateKernel(program_verify, "checkKernel", &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); - test_error_and_cleanup(errNum, CLEANUP, "error creating context\n"); - - cmd_queue3 = - clCreateCommandQueue(context2, devices[device_no], 0, &errNum); - test_error_and_cleanup(errNum, CLEANUP, - "Error: Failed to create command queue!\n"); + cmd_queue2 = clCreateCommandQueue(context, device, 0, &errNum); + test_error(errNum, "Error: Failed to create command queue!\n"); + log_info("clCreateCommandQueue successful\n"); for (int i = 0; i < 3; i++) { program_source_length = strlen(program_source_const[i]); program[i] = - clCreateProgramWithSource(context2, 1, &program_source_const[i], + clCreateProgramWithSource(context, 1, &program_source_const[i], &program_source_length, &errNum); errNum = clBuildProgram(program[i], 0, NULL, NULL, NULL, NULL); - test_error_and_cleanup(errNum, CLEANUP, - "Error: Failed to build program \n"); + test_error(errNum, "Error: Failed to build program \n"); // create the kernel - kernel2[i] = clCreateKernel(program[i], "clUpdateBuffer", &errNum); - test_error_and_cleanup(errNum, CLEANUP, "clCreateKernel failed \n"); + kernel[i] = clCreateKernel(program[i], "clUpdateBuffer", &errNum); + test_error(errNum, "clCreateKernel failed \n"); } + + program_source_const_verify = kernel_text_verify; program_source_length = strlen(program_source_const_verify); program_verify = - clCreateProgramWithSource(context2, 1, &program_source_const_verify, + clCreateProgramWithSource(context, 1, &program_source_const_verify, &program_source_length, &errNum); errNum = clBuildProgram(program_verify, 0, NULL, NULL, NULL, NULL); - test_error_and_cleanup(errNum, CLEANUP, - "Error: Failed to build program2\n"); + test_error(errNum, "Error: Failed to build program2\n"); - verify_kernel2 = clCreateKernel(program_verify, "checkKernel", &errNum); - test_error_and_cleanup(errNum, CLEANUP, "clCreateKernel failed \n"); - } + verify_kernel = clCreateKernel(program_verify, "checkKernel", &errNum); + test_error(errNum, "clCreateKernel failed \n"); - // TODO: Add support for empty list if use_fence enabled - for (VulkanExternalSemaphoreHandleType semaphoreType : - supportedSemaphoreTypes) - { - for (size_t numBuffersIdx = 0; - numBuffersIdx < ARRAY_SIZE(numBuffersList); numBuffersIdx++) + if (multiCtx) // different context guard { - uint32_t numBuffers = numBuffersList[numBuffersIdx]; - log_info("Number of buffers: %d\n", numBuffers); - for (size_t sizeIdx = 0; sizeIdx < ARRAY_SIZE(bufferSizeList); - sizeIdx++) + context2 = + clCreateContext(0, 1, &device, nullptr, nullptr, &errNum); + test_error(errNum, "error creating context\n"); + + cmd_queue3 = clCreateCommandQueue(context2, device, 0, &errNum); + test_error(errNum, "Error: Failed to create command queue!\n"); + + for (int i = 0; i < 3; i++) { - uint32_t bufferSize = bufferSizeList[sizeIdx]; - log_info( - "&&&& RUNNING vulkan_opencl_buffer test for Buffer size: " - "%d\n", - bufferSize); - if (multiImport && !multiCtx) + program_source_length = strlen(program_source_const[i]); + program[i] = clCreateProgramWithSource( + context2, 1, &program_source_const[i], + &program_source_length, &errNum); + errNum = clBuildProgram(program[i], 0, NULL, NULL, NULL, NULL); + test_error(errNum, "Error: Failed to build program \n"); + + // create the kernel + kernel2[i] = + clCreateKernel(program[i], "clUpdateBuffer", &errNum); + test_error(errNum, "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); + test_error(errNum, "Error: Failed to build program2\n"); + + verify_kernel2 = + clCreateKernel(program_verify, "checkKernel", &errNum); + test_error(errNum, "clCreateKernel failed \n"); + } + + // TODO: Add support for empty list if use_fence enabled + for (VulkanExternalSemaphoreHandleType semaphoreType : + supportedSemaphoreTypes) + { + for (size_t numBuffersIdx = 0; + numBuffersIdx < ARRAY_SIZE(numBuffersList); numBuffersIdx++) + { + 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_multi_import_same_ctx( - context, cmd_queue1, kernel, verify_kernel, vkDevice, - numBuffers, bufferSize, use_fence, semaphoreType); + 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, (cl_command_queue &)cmd_queue1, + (cl_kernel *)&kernel, (cl_kernel &)verify_kernel, + *vkDevice, numBuffers, bufferSize, use_fence, + semaphoreType); + } + else if (multiImport && multiCtx) + { + errNum = run_test_with_multi_import_diff_ctx( + context, (cl_context &)context2, + (cl_command_queue &)cmd_queue1, + (cl_command_queue &)cmd_queue3, + (cl_kernel *)&kernel, (cl_kernel *)&kernel2, + (cl_kernel &)verify_kernel, verify_kernel2, + *vkDevice, numBuffers, bufferSize, use_fence, + semaphoreType); + } + else if (numCQ == 2) + { + errNum = run_test_with_two_queue( + context, (cl_command_queue &)cmd_queue1, + (cl_command_queue &)cmd_queue2, + (cl_kernel *)&kernel, (cl_kernel &)verify_kernel, + *vkDevice, numBuffers + 1, bufferSize, use_fence, + semaphoreType); + } + else + { + errNum = run_test_with_one_queue( + context, (cl_command_queue &)cmd_queue1, + (cl_kernel *)&kernel, (cl_kernel &)verify_kernel, + *vkDevice, numBuffers, bufferSize, semaphoreType, + use_fence); + } + test_error(errNum, "func_name failed \n"); } - 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"); } } + + return errNum; } +}; -CLEANUP: - for (int i = 0; i < 3; i++) - { - if (program[i]) clReleaseProgram(program[i]); - if (kernel[i]) clReleaseKernel(kernel[i]); - } - if (cmd_queue1) clReleaseCommandQueue(cmd_queue1); - if (cmd_queue2) clReleaseCommandQueue(cmd_queue2); - if (cmd_queue3) clReleaseCommandQueue(cmd_queue3); - if (context) clReleaseContext(context); - if (context2) clReleaseContext(context2); +template struct BufferCommonBufferTest : public BufferTestBase +{ + BufferCommonBufferTest(cl_device_id device, cl_context context, + cl_command_queue queue, cl_int nelems) + : BufferTestBase(device, context, queue, nelems) + {} - if (devices) free(devices); - if (extensions) free(extensions); + cl_int Run() override { return test_buffer_common(use_fence); } +}; - return errNum; -} \ No newline at end of file +} // anonymous namespace + +int test_buffer_single_queue(cl_device_id deviceID, cl_context context, + cl_command_queue defaultQueue, int num_elements) +{ + params_reset(); + log_info("RUNNING TEST WITH ONE QUEUE...... \n\n"); + return MakeAndRunTest>( + deviceID, context, defaultQueue, num_elements); +} + +int test_buffer_multiple_queue(cl_device_id deviceID, cl_context context, + cl_command_queue defaultQueue, int num_elements) +{ + params_reset(); + numCQ = 2; + log_info("RUNNING TEST WITH TWO QUEUE...... \n\n"); + return MakeAndRunTest>( + deviceID, context, defaultQueue, num_elements); +} + +int test_buffer_multiImport_sameCtx(cl_device_id deviceID, cl_context context, + cl_command_queue defaultQueue, + int num_elements) +{ + params_reset(); + multiImport = true; + log_info("RUNNING TEST WITH MULTIPLE DEVICE MEMORY IMPORT " + "IN SAME CONTEXT...... \n\n"); + return MakeAndRunTest>( + deviceID, context, defaultQueue, num_elements); +} +int test_buffer_multiImport_diffCtx(cl_device_id deviceID, cl_context context, + cl_command_queue defaultQueue, + int num_elements) +{ + params_reset(); + multiImport = true; + multiCtx = true; + log_info("RUNNING TEST WITH MULTIPLE DEVICE MEMORY IMPORT " + "IN DIFFERENT CONTEXT...... \n\n"); + return MakeAndRunTest>( + deviceID, context, defaultQueue, num_elements); +} +int test_buffer_single_queue_fence(cl_device_id deviceID, cl_context context, + cl_command_queue defaultQueue, + int num_elements) +{ + params_reset(); + log_info("RUNNING TEST WITH ONE QUEUE...... \n\n"); + + return MakeAndRunTest>( + deviceID, context, defaultQueue, num_elements); +} + +int test_buffer_multiple_queue_fence(cl_device_id deviceID, cl_context context, + cl_command_queue defaultQueue, + int num_elements) +{ + params_reset(); + numCQ = 2; + log_info("RUNNING TEST WITH TWO QUEUE...... \n\n"); + return MakeAndRunTest>( + deviceID, context, defaultQueue, num_elements); +} + +int test_buffer_multiImport_sameCtx_fence(cl_device_id deviceID, + cl_context context, + cl_command_queue defaultQueue, + int num_elements) +{ + params_reset(); + multiImport = true; + log_info("RUNNING TEST WITH MULTIPLE DEVICE MEMORY IMPORT " + "IN SAME CONTEXT...... \n\n"); + return MakeAndRunTest>( + deviceID, context, defaultQueue, num_elements); +} + +int test_buffer_multiImport_diffCtx_fence(cl_device_id deviceID, + cl_context context, + cl_command_queue defaultQueue, + int num_elements) +{ + params_reset(); + multiImport = true; + multiCtx = true; + log_info("RUNNING TEST WITH MULTIPLE DEVICE MEMORY IMPORT " + "IN DIFFERENT CONTEXT...... \n\n"); + return MakeAndRunTest>( + deviceID, context, defaultQueue, num_elements); +} diff --git a/test_conformance/vulkan/test_vulkan_interop_image.cpp b/test_conformance/vulkan/test_vulkan_interop_image.cpp index a3c8de99..7808ef64 100644 --- a/test_conformance/vulkan/test_vulkan_interop_image.cpp +++ b/test_conformance/vulkan/test_vulkan_interop_image.cpp @@ -1,5 +1,5 @@ // -// Copyright (c) 2022 The Khronos Group Inc. +// Copyright (c) 2024 The Khronos Group Inc. // // Licensed under the Apache License, Version 2.0 (the "License"); // you may not use this file except in compliance with the License. @@ -19,7 +19,11 @@ #include "harness/errorHelpers.h" #include "harness/os_helpers.h" #include -#include "deviceInfo.h" + +#include "vulkan_test_base.h" +#include "opencl_vulkan_wrapper.hpp" + +namespace { #define MAX_2D_IMAGES 5 #define MAX_2D_IMAGE_WIDTH 1024 @@ -46,14 +50,13 @@ ASSERT(0); \ } -namespace { struct Params { uint32_t numImage2DDescriptors; }; -} -static cl_uchar uuid[CL_UUID_SIZE_KHR]; -static cl_device_id deviceId = NULL; + +cl_uchar uuid[CL_UUID_SIZE_KHR]; +cl_device_id deviceId = NULL; size_t max_width = MAX_2D_IMAGE_WIDTH; size_t max_height = MAX_2D_IMAGE_HEIGHT; @@ -245,7 +248,7 @@ int run_test_with_two_queue( VulkanCommandPool vkCommandPool(vkDevice); VulkanCommandBuffer vkCopyCommandBuffer(vkDevice, vkCommandPool); VulkanCommandBuffer vkShaderCommandBuffer(vkDevice, vkCommandPool); - VulkanQueue &vkQueue = vkDevice.getQueue(); + VulkanQueue &vkQueue = vkDevice.getQueue(getVulkanQueueFamily()); VulkanSemaphore vkVk2CLSemaphore(vkDevice, vkExternalSemaphoreHandleType); VulkanSemaphore vkCl2VkSemaphore(vkDevice, vkExternalSemaphoreHandleType); @@ -857,7 +860,7 @@ int run_test_with_one_queue( VulkanCommandPool vkCommandPool(vkDevice); VulkanCommandBuffer vkCopyCommandBuffer(vkDevice, vkCommandPool); VulkanCommandBuffer vkShaderCommandBuffer(vkDevice, vkCommandPool); - VulkanQueue &vkQueue = vkDevice.getQueue(); + VulkanQueue &vkQueue = vkDevice.getQueue(getVulkanQueueFamily()); VulkanSemaphore vkVk2CLSemaphore(vkDevice, vkExternalSemaphoreHandleType); VulkanSemaphore vkCl2VkSemaphore(vkDevice, vkExternalSemaphoreHandleType); @@ -1352,262 +1355,185 @@ CLEANUP: return err; } -int test_image_common(cl_device_id device_, cl_context context_, - cl_command_queue queue_, int numElements_) +struct ImageCommonTest : public VulkanTestBase { - int current_device = 0; - int device_count = 0; - int devices_prohibited = 0; - cl_int err = CL_SUCCESS; - cl_platform_id platform = NULL; - size_t extensionSize = 0; - cl_uint num_devices = 0; - cl_uint device_no = 0; - cl_device_id *devices; - char *extensions = NULL; - const char *program_source_const; - cl_command_queue cmd_queue1 = NULL; - cl_command_queue cmd_queue2 = NULL; - cl_context context = NULL; - const uint32_t num_kernels = ARRAY_SIZE(num2DImagesList) + 1; - // One kernel for Cross-CQ case - const uint32_t num_kernel_types = 3; - const char *kernel_source[num_kernels] = { kernel_text_numImage_1, - kernel_text_numImage_2, - kernel_text_numImage_4 }; - char source_1[4096]; - char source_2[4096]; - char source_3[4096]; - size_t program_source_length; - 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; + ImageCommonTest(cl_device_id device, cl_context context, + cl_command_queue queue, cl_int nelems) + : VulkanTestBase(device, context, queue, nelems) + {} - VulkanDevice vkDevice; - - cl_context_properties contextProperties[] = { CL_CONTEXT_PLATFORM, 0, 0 }; - // get the platform ID - err = clGetPlatformIDs(1, &platform, NULL); - test_error_and_cleanup(err, CLEANUP, "Error: Failed to get platform\n"); - - err = clGetDeviceIDs(platform, CL_DEVICE_TYPE_GPU, 0, NULL, &num_devices); - 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) + int test_image_common() { - test_fail_and_cleanup(err, CLEANUP, - "Unable to allocate memory for devices\n"); - } - err = clGetDeviceIDs(platform, CL_DEVICE_TYPE_GPU, num_devices, devices, - NULL); - 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++) - { - err = clGetDeviceInfo(devices[device_no], CL_DEVICE_EXTENSIONS, 0, NULL, - &extensionSize); - if (CL_SUCCESS != err) - { - print_error( - err, - "Error in clGetDeviceInfo for getting device_extension size\n"); - goto CLEANUP; - } - extensions = (char *)malloc(extensionSize); - if (NULL == extensions) - { - err = CL_OUT_OF_HOST_MEMORY; - print_error(err, "Unable to allocate memory for extensions\n"); - goto CLEANUP; - } - err = clGetDeviceInfo(devices[device_no], CL_DEVICE_EXTENSIONS, - extensionSize, extensions, NULL); - if (CL_SUCCESS != err) - { - print_error( - err, "Error in clGetDeviceInfo for getting device_extension\n"); - goto CLEANUP; - } - err = clGetDeviceInfo(devices[device_no], CL_DEVICE_UUID_KHR, - CL_UUID_SIZE_KHR, uuid, NULL); - test_error_and_cleanup(err, CLEANUP, - "clGetDeviceInfo failed with error"); + cl_int err = CL_SUCCESS; + clCommandQueueWrapper cmd_queue1; + clCommandQueueWrapper cmd_queue2; + const uint32_t num_kernels = ARRAY_SIZE(num2DImagesList) + 1; + // One kernel for Cross-CQ case + const uint32_t num_kernel_types = 3; + const char *kernel_source[num_kernels] = { kernel_text_numImage_1, + kernel_text_numImage_2, + kernel_text_numImage_4 }; + char source_1[4096]; + char source_2[4096]; + char source_3[4096]; + size_t program_source_length; + clProgramWrapper program[num_kernel_types] = { NULL }; + clKernelWrapper kernel_float[num_kernels] = { NULL }; + clKernelWrapper kernel_signed[num_kernels] = { NULL }; + clKernelWrapper kernel_unsigned[num_kernels] = { NULL }; + clMemWrapper external_mem_image1; + clMemWrapper external_mem_image2; + std::vector supportedSemaphoreTypes; supportedSemaphoreTypes = - getSupportedInteropExternalSemaphoreHandleTypes(devices[device_no], - vkDevice); + getSupportedInteropExternalSemaphoreHandleTypes(device, *vkDevice); // If device does not support any semaphores, try the next one if (supportedSemaphoreTypes.empty()) { - continue; + log_info("Device does not support any semaphores!\n"); + return TEST_SKIPPED_ITSELF; } - err = - memcmp(uuid, vkDevice.getPhysicalDevice().getUUID(), VK_UUID_SIZE); - if (err == 0) + deviceId = device; + + err = setMaxImageDimensions(deviceId, max_width, max_height); + test_error(err, "error setting max image dimensions"); + + log_info("Set max_width to %zu and max_height to %zu\n", max_width, + max_height); + + log_info("Successfully created context !!!\n"); + + cmd_queue1 = clCreateCommandQueue(context, deviceId, 0, &err); + test_error(err, "Error: Failed to create command queue!\n"); + + log_info("clCreateCommandQueue successfull \n"); + + cmd_queue2 = clCreateCommandQueue(context, deviceId, 0, &err); + test_error(err, "Error: Failed to create command queue!\n"); + + log_info("clCreateCommandQueue2 successful \n"); + + for (int i = 0; i < num_kernels; i++) { - break; + switch (i) + { + case 0: + sprintf(source_1, kernel_source[i], "float4", "f", "float4", + "f", "f", "f"); + sprintf(source_2, kernel_source[i], "int4", "i", "int4", + "i", "i", "i"); + sprintf(source_3, kernel_source[i], "uint4", "ui", "uint4", + "ui", "ui", "ui"); + break; + case 1: + sprintf(source_1, kernel_source[i], "float4", "f", "float4", + "f", "float4", "f", "float4", "f", "f", "f", "f", + "f"); + sprintf(source_2, kernel_source[i], "int4", "i", "int4", + "i", "int4", "i", "int4", "i", "i", "i", "i", "i"); + sprintf(source_3, kernel_source[i], "uint4", "ui", "uint4", + "ui", "uint4", "ui", "uint4", "ui", "ui", "ui", + "ui", "ui"); + break; + case 2: + sprintf(source_1, kernel_source[i], "float4", "f", "float4", + "f", "float4", "f", "float4", "f", "float4", "f", + "float4", "f", "float4", "f", "float4", "f", "f", + "f", "f", "f", "f", "f", "f", "f"); + sprintf(source_2, kernel_source[i], "int4", "i", "int4", + "i", "int4", "i", "int4", "i", "int4", "i", "int4", + "i", "int4", "i", "int4", "i", "i", "i", "i", "i", + "i", "i", "i", "i"); + sprintf(source_3, kernel_source[i], "uint4", "ui", "uint4", + "ui", "uint4", "ui", "uint4", "ui", "uint4", "ui", + "uint4", "ui", "uint4", "ui", "uint4", "ui", "ui", + "ui", "ui", "ui", "ui", "ui", "ui", "ui"); + break; + case 3: + // Addtional case for creating updateKernelCQ2 which takes + // two images + sprintf(source_1, kernel_source[1], "float4", "f", "float4", + "f", "float4", "f", "float4", "f", "f", "f", "f", + "f"); + sprintf(source_2, kernel_source[1], "int4", "i", "int4", + "i", "int4", "i", "int4", "i", "i", "i", "i", "i"); + sprintf(source_3, kernel_source[1], "uint4", "ui", "uint4", + "ui", "uint4", "ui", "uint4", "ui", "ui", "ui", + "ui", "ui"); + break; + } + const char *sourceTexts[num_kernel_types] = { source_1, source_2, + source_3 }; + for (int k = 0; k < num_kernel_types; k++) + { + program_source_length = strlen(sourceTexts[k]); + program[k] = clCreateProgramWithSource( + context, 1, &sourceTexts[k], &program_source_length, &err); + err |= clBuildProgram(program[k], 0, NULL, NULL, NULL, NULL); + } + test_error(err, "Error: Failed to build program"); + + // create the kernel + kernel_float[i] = clCreateKernel(program[0], "image2DKernel", &err); + test_error(err, "clCreateKernel failed"); + + kernel_signed[i] = + clCreateKernel(program[1], "image2DKernel", &err); + test_error(err, "clCreateKernel failed"); + + kernel_unsigned[i] = + clCreateKernel(program[2], "image2DKernel", &err); + test_error(err, "clCreateKernel failed "); } + for (VulkanExternalSemaphoreHandleType externalSemaphoreType : + supportedSemaphoreTypes) + { + if (numCQ == 2) + { + err = run_test_with_two_queue( + context, (cl_command_queue &)cmd_queue1, + (cl_command_queue &)cmd_queue2, + (cl_kernel *)kernel_unsigned, (cl_kernel *)kernel_signed, + (cl_kernel *)kernel_float, *vkDevice, + externalSemaphoreType); + } + else + { + err = run_test_with_one_queue( + context, (cl_command_queue &)cmd_queue1, + (cl_kernel *)kernel_unsigned, (cl_kernel *)kernel_signed, + (cl_kernel *)kernel_float, *vkDevice, + externalSemaphoreType); + } + test_error(err, "func_name failed \n"); + } + + return err; } - if (supportedSemaphoreTypes.empty()) - { - test_fail_and_cleanup( - err, CLEANUP, "No devices found that support OpenCL semaphores\n"); - } + cl_int Run() override { return test_image_common(); } +}; - if (device_no >= num_devices) - { - 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); - test_error_and_cleanup(err, CLEANUP, "error setting max image dimensions"); +} // anonymous namespace - log_info("Set max_width to %zu and max_height to %zu\n", max_width, - max_height); - context = clCreateContextFromType(contextProperties, CL_DEVICE_TYPE_GPU, - NULL, NULL, &err); - test_error_and_cleanup(err, CLEANUP, "error creating context"); +int test_image_single_queue(cl_device_id deviceID, cl_context context, + cl_command_queue defaultQueue, int num_elements) +{ + params_reset(); + log_info("RUNNING TEST WITH ONE QUEUE...... \n\n"); - log_info("Successfully created context !!!\n"); + return MakeAndRunTest(deviceID, context, defaultQueue, + num_elements); +} - cmd_queue1 = clCreateCommandQueue(context, devices[device_no], 0, &err); - 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); - 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++) - { - switch (i) - { - case 0: - sprintf(source_1, kernel_source[i], "float4", "f", "float4", - "f", "f", "f"); - sprintf(source_2, kernel_source[i], "int4", "i", "int4", "i", - "i", "i"); - sprintf(source_3, kernel_source[i], "uint4", "ui", "uint4", - "ui", "ui", "ui"); - break; - case 1: - sprintf(source_1, kernel_source[i], "float4", "f", "float4", - "f", "float4", "f", "float4", "f", "f", "f", "f", "f"); - sprintf(source_2, kernel_source[i], "int4", "i", "int4", "i", - "int4", "i", "int4", "i", "i", "i", "i", "i"); - sprintf(source_3, kernel_source[i], "uint4", "ui", "uint4", - "ui", "uint4", "ui", "uint4", "ui", "ui", "ui", "ui", - "ui"); - break; - case 2: - sprintf(source_1, kernel_source[i], "float4", "f", "float4", - "f", "float4", "f", "float4", "f", "float4", "f", - "float4", "f", "float4", "f", "float4", "f", "f", "f", - "f", "f", "f", "f", "f", "f"); - sprintf(source_2, kernel_source[i], "int4", "i", "int4", "i", - "int4", "i", "int4", "i", "int4", "i", "int4", "i", - "int4", "i", "int4", "i", "i", "i", "i", "i", "i", "i", - "i", "i"); - sprintf(source_3, kernel_source[i], "uint4", "ui", "uint4", - "ui", "uint4", "ui", "uint4", "ui", "uint4", "ui", - "uint4", "ui", "uint4", "ui", "uint4", "ui", "ui", "ui", - "ui", "ui", "ui", "ui", "ui", "ui"); - break; - case 3: - // Addtional case for creating updateKernelCQ2 which takes two - // images - sprintf(source_1, kernel_source[1], "float4", "f", "float4", - "f", "float4", "f", "float4", "f", "f", "f", "f", "f"); - sprintf(source_2, kernel_source[1], "int4", "i", "int4", "i", - "int4", "i", "int4", "i", "i", "i", "i", "i"); - sprintf(source_3, kernel_source[1], "uint4", "ui", "uint4", - "ui", "uint4", "ui", "uint4", "ui", "ui", "ui", "ui", - "ui"); - break; - } - const char *sourceTexts[num_kernel_types] = { source_1, source_2, - source_3 }; - for (int k = 0; k < num_kernel_types; k++) - { - program_source_length = strlen(sourceTexts[k]); - program[k] = clCreateProgramWithSource( - 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"); - - // create the kernel - kernel_float[i] = clCreateKernel(program[0], "image2DKernel", &err); - test_error_and_cleanup(err, CLEANUP, "clCreateKernel failed"); - - kernel_signed[i] = clCreateKernel(program[1], "image2DKernel", &err); - 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) - { - err = run_test_with_two_queue( - context, cmd_queue1, cmd_queue2, kernel_unsigned, kernel_signed, - kernel_float, vkDevice, externalSemaphoreType); - } - else - { - err = run_test_with_one_queue(context, cmd_queue1, kernel_unsigned, - kernel_signed, kernel_float, vkDevice, - externalSemaphoreType); - } - } -CLEANUP: - for (int i = 0; i < num_kernels; i++) - { - if (kernel_float[i]) - { - clReleaseKernel(kernel_float[i]); - } - if (kernel_unsigned[i]) - { - clReleaseKernel(kernel_unsigned[i]); - } - if (kernel_signed[i]) - { - clReleaseKernel(kernel_signed[i]); - } - } - for (int i = 0; i < num_kernel_types; i++) - { - if (program[i]) - { - clReleaseProgram(program[i]); - } - } - if (cmd_queue1) clReleaseCommandQueue(cmd_queue1); - if (cmd_queue2) clReleaseCommandQueue(cmd_queue2); - if (context) clReleaseContext(context); - - if (extensions) free(extensions); - if (devices) free(devices); - - return err; -} \ No newline at end of file +int test_image_multiple_queue(cl_device_id deviceID, cl_context context, + cl_command_queue defaultQueue, int num_elements) +{ + params_reset(); + numCQ = 2; + log_info("RUNNING TEST WITH TWO QUEUE...... \n\n"); + return MakeAndRunTest(deviceID, context, defaultQueue, + num_elements); +} diff --git a/test_conformance/vulkan/test_vulkan_platform_device_info.cpp b/test_conformance/vulkan/test_vulkan_platform_device_info.cpp index 1c25c0f5..eaf963c9 100644 --- a/test_conformance/vulkan/test_vulkan_platform_device_info.cpp +++ b/test_conformance/vulkan/test_vulkan_platform_device_info.cpp @@ -1,5 +1,5 @@ // -// Copyright (c) 2022 The Khronos Group Inc. +// Copyright (c) 2024 The Khronos Group Inc. // // Licensed under the Apache License, Version 2.0 (the "License"); // you may not use this file except in compliance with the License. @@ -22,6 +22,10 @@ #include #include +#include "vulkan_test_base.h" + +namespace { + typedef struct { cl_uint info; @@ -29,183 +33,216 @@ typedef struct } _info; _info platform_info_table[] = { -#define STRING(x) \ +#define PLATFORM_INFO_STRING(x) \ { \ x, #x \ } - STRING(CL_PLATFORM_EXTERNAL_MEMORY_IMPORT_HANDLE_TYPES_KHR), - STRING(CL_PLATFORM_SEMAPHORE_EXPORT_HANDLE_TYPES_KHR), - STRING(CL_PLATFORM_SEMAPHORE_IMPORT_HANDLE_TYPES_KHR) -#undef STRING + PLATFORM_INFO_STRING(CL_PLATFORM_EXTERNAL_MEMORY_IMPORT_HANDLE_TYPES_KHR), + PLATFORM_INFO_STRING(CL_PLATFORM_SEMAPHORE_EXPORT_HANDLE_TYPES_KHR), + PLATFORM_INFO_STRING(CL_PLATFORM_SEMAPHORE_IMPORT_HANDLE_TYPES_KHR) +#undef PLATFORM_INFO_STRING }; _info device_info_table[] = { -#define STRING(x) \ +#define DEVICE_INFO_STRING(x) \ { \ x, #x \ } - STRING(CL_DEVICE_SEMAPHORE_IMPORT_HANDLE_TYPES_KHR), - STRING(CL_DEVICE_SEMAPHORE_EXPORT_HANDLE_TYPES_KHR), - STRING(CL_DEVICE_EXTERNAL_MEMORY_IMPORT_HANDLE_TYPES_KHR) -#undef STRING + DEVICE_INFO_STRING(CL_DEVICE_SEMAPHORE_IMPORT_HANDLE_TYPES_KHR), + DEVICE_INFO_STRING(CL_DEVICE_SEMAPHORE_EXPORT_HANDLE_TYPES_KHR), + DEVICE_INFO_STRING(CL_DEVICE_EXTERNAL_MEMORY_IMPORT_HANDLE_TYPES_KHR) +#undef DEVICE_INFO_STRING }; -int test_platform_info(cl_device_id deviceID, cl_context _context, - cl_command_queue _queue, int num_elements) +struct PlatformInfoTest : public VulkanTestBase { - cl_uint i; - cl_platform_id platform = getPlatformFromDevice(deviceID); - cl_int errNum; - cl_uint *handle_type; - size_t handle_type_size = 0; - cl_uint num_handles = 0; - cl_bool external_mem_extn_available = - is_platform_extension_available(platform, "cl_khr_external_semaphore"); - cl_bool external_sema_extn_available = - is_platform_extension_available(platform, "cl_khr_external_memory"); - cl_bool supports_atleast_one_sema_query = false; + PlatformInfoTest(cl_device_id device, cl_context context, + cl_command_queue queue, cl_int nelems) + : VulkanTestBase(device, context, queue, nelems) + {} - if (!external_mem_extn_available && !external_sema_extn_available) + cl_int Run() override { - log_info("Platform does not support 'cl_khr_external_semaphore' " - "and 'cl_khr_external_memory'. Skipping the test.\n"); - return TEST_SKIPPED_ITSELF; - } + cl_uint i; + cl_platform_id platform = getPlatformFromDevice(device); + cl_int errNum; + cl_uint *handle_type; + size_t handle_type_size = 0; + cl_uint num_handles = 0; + cl_bool external_mem_extn_available = is_platform_extension_available( + platform, "cl_khr_external_semaphore"); + cl_bool external_sema_extn_available = + is_platform_extension_available(platform, "cl_khr_external_memory"); + cl_bool supports_atleast_one_sema_query = false; - log_info("Platform (id %lu) info:\n", (unsigned long)platform); - - for (i = 0; - i < sizeof(platform_info_table) / sizeof(platform_info_table[0]); i++) - { - errNum = clGetPlatformInfo(platform, platform_info_table[i].info, 0, - NULL, &handle_type_size); - test_error(errNum, "clGetPlatformInfo failed"); - - if (handle_type_size == 0) + if (!external_mem_extn_available && !external_sema_extn_available) { - if (platform_info_table[i].info - == CL_PLATFORM_EXTERNAL_MEMORY_IMPORT_HANDLE_TYPES_KHR - && external_mem_extn_available) + log_info("Platform does not support 'cl_khr_external_semaphore' " + "and 'cl_khr_external_memory'. Skipping the test.\n"); + return TEST_SKIPPED_ITSELF; + } + + log_info("Platform (id %lu) info:\n", (unsigned long)platform); + + for (i = 0; + i < sizeof(platform_info_table) / sizeof(platform_info_table[0]); + i++) + { + errNum = clGetPlatformInfo(platform, platform_info_table[i].info, 0, + NULL, &handle_type_size); + test_error(errNum, "clGetPlatformInfo failed"); + + if (handle_type_size == 0) { - test_fail( - "External memory import handle types should be reported if " - "cl_khr_external_memory is available.\n"); + if (platform_info_table[i].info + == CL_PLATFORM_EXTERNAL_MEMORY_IMPORT_HANDLE_TYPES_KHR + && external_mem_extn_available) + { + test_fail("External memory import handle types should be " + "reported if " + "cl_khr_external_memory is available.\n"); + } + log_info("%s not supported. Skipping the query.\n", + platform_info_table[i].name); + continue; + } + + if ((platform_info_table[i].info + == CL_PLATFORM_SEMAPHORE_EXPORT_HANDLE_TYPES_KHR) + || (platform_info_table[i].info + == CL_PLATFORM_SEMAPHORE_IMPORT_HANDLE_TYPES_KHR)) + { + supports_atleast_one_sema_query = true; + } + + num_handles = handle_type_size / sizeof(cl_uint); + handle_type = (cl_uint *)malloc(handle_type_size); + errNum = clGetPlatformInfo(platform, platform_info_table[i].info, + handle_type_size, handle_type, NULL); + test_error(errNum, "clGetPlatformInfo failed"); + + log_info("%s: \n", platform_info_table[i].name); + while (num_handles--) + { + log_info("%x \n", handle_type[num_handles]); + } + if (handle_type) + { + free(handle_type); } - log_info("%s not supported. Skipping the query.\n", - platform_info_table[i].name); - continue; } - if ((platform_info_table[i].info - == CL_PLATFORM_SEMAPHORE_EXPORT_HANDLE_TYPES_KHR) - || (platform_info_table[i].info - == CL_PLATFORM_SEMAPHORE_IMPORT_HANDLE_TYPES_KHR)) + if (external_sema_extn_available && !supports_atleast_one_sema_query) { - supports_atleast_one_sema_query = true; + log_info( + "External semaphore import/export or both should be supported " + "if cl_khr_external_semaphore is available.\n"); + return TEST_FAIL; } - num_handles = handle_type_size / sizeof(cl_uint); - handle_type = (cl_uint *)malloc(handle_type_size); - errNum = clGetPlatformInfo(platform, platform_info_table[i].info, - handle_type_size, handle_type, NULL); - test_error(errNum, "clGetPlatformInfo failed"); - - log_info("%s: \n", platform_info_table[i].name); - while (num_handles--) - { - log_info("%x \n", handle_type[num_handles]); - } - if (handle_type) - { - free(handle_type); - } + return TEST_PASS; } +}; - if (external_sema_extn_available && !supports_atleast_one_sema_query) +struct DeviceInfoTest : public VulkanTestBase +{ + DeviceInfoTest(cl_device_id device, cl_context context, + cl_command_queue queue, cl_int nelems) + : VulkanTestBase(device, context, queue, nelems) + {} + + cl_int Run() override { - log_info("External semaphore import/export or both should be supported " - "if cl_khr_external_semaphore is available.\n"); - return TEST_FAIL; - } + cl_uint j; + cl_uint *handle_type; + size_t handle_type_size = 0; + cl_uint num_handles = 0; + cl_int errNum = CL_SUCCESS; + cl_bool external_mem_extn_available = + is_extension_available(device, "cl_khr_external_memory"); + cl_bool external_sema_extn_available = + is_extension_available(device, "cl_khr_external_semaphore"); + cl_bool supports_atleast_one_sema_query = false; - return TEST_PASS; + if (!external_mem_extn_available && !external_sema_extn_available) + { + log_info("Device does not support 'cl_khr_external_semaphore' " + "and 'cl_khr_external_memory'. Skipping the test.\n"); + return TEST_SKIPPED_ITSELF; + } + + for (j = 0; + j < sizeof(device_info_table) / sizeof(device_info_table[0]); j++) + { + errNum = clGetDeviceInfo(device, device_info_table[j].info, 0, NULL, + &handle_type_size); + test_error(errNum, "clGetDeviceInfo failed"); + + if (handle_type_size == 0) + { + if (device_info_table[j].info + == CL_DEVICE_EXTERNAL_MEMORY_IMPORT_HANDLE_TYPES_KHR + && external_mem_extn_available) + { + test_fail("External memory import handle types should be " + "reported if " + "cl_khr_external_memory is available.\n"); + } + log_info("%s not supported. Skipping the query.\n", + device_info_table[j].name); + continue; + } + + if ((device_info_table[j].info + == CL_DEVICE_SEMAPHORE_EXPORT_HANDLE_TYPES_KHR) + || (device_info_table[j].info + == CL_DEVICE_SEMAPHORE_IMPORT_HANDLE_TYPES_KHR)) + { + supports_atleast_one_sema_query = true; + } + + num_handles = handle_type_size / sizeof(cl_uint); + handle_type = (cl_uint *)malloc(handle_type_size); + + errNum = clGetDeviceInfo(device, device_info_table[j].info, + handle_type_size, handle_type, NULL); + test_error(errNum, "clGetDeviceInfo failed"); + + log_info("%s: \n", device_info_table[j].name); + while (num_handles--) + { + log_info("%x \n", handle_type[num_handles]); + } + if (handle_type) + { + free(handle_type); + } + } + + if (external_sema_extn_available && !supports_atleast_one_sema_query) + { + log_info( + "External semaphore import/export or both should be supported " + "if cl_khr_external_semaphore is available.\n"); + return TEST_FAIL; + } + + return TEST_PASS; + } +}; + +} // anonymous namespace + +int test_platform_info(cl_device_id deviceID, cl_context context, + cl_command_queue defaultQueue, int num_elements) +{ + return MakeAndRunTest(deviceID, context, defaultQueue, + num_elements); } -int test_device_info(cl_device_id deviceID, cl_context _context, - cl_command_queue _queue, int num_elements) +int test_device_info(cl_device_id deviceID, cl_context context, + cl_command_queue defaultQueue, int num_elements) { - cl_uint j; - cl_uint *handle_type; - size_t handle_type_size = 0; - cl_uint num_handles = 0; - cl_int errNum = CL_SUCCESS; - cl_bool external_mem_extn_available = - is_extension_available(deviceID, "cl_khr_external_memory"); - cl_bool external_sema_extn_available = - is_extension_available(deviceID, "cl_khr_external_semaphore"); - cl_bool supports_atleast_one_sema_query = false; - - if (!external_mem_extn_available && !external_sema_extn_available) - { - log_info("Device does not support 'cl_khr_external_semaphore' " - "and 'cl_khr_external_memory'. Skipping the test.\n"); - return TEST_SKIPPED_ITSELF; - } - - for (j = 0; j < sizeof(device_info_table) / sizeof(device_info_table[0]); - j++) - { - errNum = clGetDeviceInfo(deviceID, device_info_table[j].info, 0, NULL, - &handle_type_size); - test_error(errNum, "clGetDeviceInfo failed"); - - if (handle_type_size == 0) - { - if (device_info_table[j].info - == CL_DEVICE_EXTERNAL_MEMORY_IMPORT_HANDLE_TYPES_KHR - && external_mem_extn_available) - { - test_fail( - "External memory import handle types should be reported if " - "cl_khr_external_memory is available.\n"); - } - log_info("%s not supported. Skipping the query.\n", - device_info_table[j].name); - continue; - } - - if ((device_info_table[j].info - == CL_DEVICE_SEMAPHORE_EXPORT_HANDLE_TYPES_KHR) - || (device_info_table[j].info - == CL_DEVICE_SEMAPHORE_IMPORT_HANDLE_TYPES_KHR)) - { - supports_atleast_one_sema_query = true; - } - - num_handles = handle_type_size / sizeof(cl_uint); - handle_type = (cl_uint *)malloc(handle_type_size); - - errNum = clGetDeviceInfo(deviceID, device_info_table[j].info, - handle_type_size, handle_type, NULL); - test_error(errNum, "clGetDeviceInfo failed"); - - log_info("%s: \n", device_info_table[j].name); - while (num_handles--) - { - log_info("%x \n", handle_type[num_handles]); - } - if (handle_type) - { - free(handle_type); - } - } - - if (external_sema_extn_available && !supports_atleast_one_sema_query) - { - log_info("External semaphore import/export or both should be supported " - "if cl_khr_external_semaphore is available.\n"); - return TEST_FAIL; - } - - return TEST_PASS; + return MakeAndRunTest(deviceID, context, defaultQueue, + num_elements); } diff --git a/test_conformance/vulkan/vulkan_test_base.h b/test_conformance/vulkan/vulkan_test_base.h new file mode 100644 index 00000000..d4cfa684 --- /dev/null +++ b/test_conformance/vulkan/vulkan_test_base.h @@ -0,0 +1,129 @@ +// +// Copyright (c) 2024 The Khronos Group Inc. +// +// Licensed under the Apache License, Version 2.0 (the "License"); +// you may not use this file except in compliance with the License. +// You may obtain a copy of the License at +// +// http://www.apache.org/licenses/LICENSE-2.0 +// +// Unless required by applicable law or agreed to in writing, software +// distributed under the License is distributed on an "AS IS" BASIS, +// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +// See the License for the specific language governing permissions and +// limitations under the License. + +#ifndef CL_VULKAN_TEST_BASE_H +#define CL_VULKAN_TEST_BASE_H + +#include + +#include +#include + +#include "vulkan_interop_common.hpp" + +#include "harness/deviceInfo.h" +#include "harness/testHarness.h" +#include "harness/typeWrappers.h" + +inline void params_reset() +{ + numCQ = 1; + multiImport = false; + multiCtx = false; +} + +struct VulkanTestBase +{ + VulkanTestBase(cl_device_id device, cl_context context, + cl_command_queue queue, cl_int nelems) + : device(device), context(context), num_elems(nelems) + { + vkDevice.reset( + new VulkanDevice(getAssociatedVulkanPhysicalDevice(device))); + + if (!(is_extension_available(device, "cl_khr_external_memory") + && is_extension_available(device, "cl_khr_external_semaphore"))) + { + log_info("Device does not support cl_khr_external_memory " + "or cl_khr_external_semaphore\n"); + log_info(" TEST SKIPPED\n"); + throw std::runtime_error("VulkanTestBase not supported"); + } + + cl_platform_id platform; + cl_int error = clGetDeviceInfo(device, CL_DEVICE_PLATFORM, + sizeof(cl_platform_id), &platform, NULL); + if (error != CL_SUCCESS) + throw std::runtime_error( + "clGetDeviceInfo for CL_DEVICE_PLATFORM failed"); + + + // verify whether selected device is one of the type CL_DEVICE_TYPE_GPU + cl_uint num_devices = 0; + error = + clGetDeviceIDs(platform, CL_DEVICE_TYPE_GPU, 0, NULL, &num_devices); + if (CL_SUCCESS != error) + throw std::runtime_error( + "clGetDeviceIDs failed in returning of devices"); + + std::vector devices(num_devices); + error = clGetDeviceIDs(platform, CL_DEVICE_TYPE_GPU, num_devices, + devices.data(), NULL); + + bool found_gpu_match = false; + for (cl_uint i = 0; i < num_devices; i++) + if (devices[i] == device) + { + found_gpu_match = true; + break; + } + + if (!found_gpu_match) + throw std::runtime_error( + "Vulkan tests can only run on a GPU device."); + + init_cl_vk_ext(platform, 1, &device); + } + + virtual cl_int Run() = 0; + +protected: + cl_device_id device = nullptr; + cl_context context = nullptr; + clCommandQueueWrapper queue = nullptr; + cl_int num_elems = 0; + std::unique_ptr vkDevice; +}; + +template +int MakeAndRunTest(cl_device_id device, cl_context context, + cl_command_queue queue, cl_int nelems) +{ + if (!checkVkSupport()) + { + log_info("Vulkan supported GPU not found \n"); + log_info("TEST SKIPPED \n"); + return TEST_SKIPPED_ITSELF; + } + + cl_int status = TEST_PASS; + try + { + // moved from original test - do we want to stick to that ? + cl_int numElementsToUse = 1024; + + auto test_fixture = + T(device, context, queue, /*nelems*/ numElementsToUse); + status = test_fixture.Run(); + } catch (const std::runtime_error &e) + { + log_error("%s", e.what()); + return TEST_FAIL; + } + + return status; +} + +#endif // CL_VULKAN_TEST_BASE_H