diff --git a/test_conformance/common/vulkan_wrapper/vulkan_wrapper.cpp b/test_conformance/common/vulkan_wrapper/vulkan_wrapper.cpp index 3ce4af6b0..21d8f2260 100644 --- a/test_conformance/common/vulkan_wrapper/vulkan_wrapper.cpp +++ b/test_conformance/common/vulkan_wrapper/vulkan_wrapper.cpp @@ -604,6 +604,37 @@ VulkanQueue &VulkanDevice::getQueue(const VulkanQueueFamily &queueFamily, VulkanDevice::operator VkDevice() const { return m_vkDevice; } +//////////////////////////////// +// VulkanFence implementation // +//////////////////////////////// + +VulkanFence::VulkanFence(const VulkanDevice &vkDevice) +{ + + device = vkDevice; + + VkFenceCreateInfo fenceInfo{}; + fenceInfo.sType = VK_STRUCTURE_TYPE_FENCE_CREATE_INFO; + fenceInfo.pNext = nullptr; + fenceInfo.flags = 0; + + VkResult vkStatus = vkCreateFence(device, &fenceInfo, nullptr, &fence); + + if (vkStatus != VK_SUCCESS) + { + throw std::runtime_error("Error: Failed create fence."); + } +} + +VulkanFence::~VulkanFence() { vkDestroyFence(device, fence, nullptr); } + +void VulkanFence::reset() { vkResetFences(device, 1, &fence); } + +void VulkanFence::wait() +{ + vkWaitForFences(device, 1, &fence, VK_TRUE, UINT64_MAX); +} + //////////////////////////////// // VulkanQueue implementation // //////////////////////////////// @@ -615,6 +646,22 @@ VulkanQueue::VulkanQueue(VkQueue vkQueue): m_vkQueue(vkQueue) {} VulkanQueue::~VulkanQueue() {} +void VulkanQueue::submit(const VulkanCommandBuffer &commandBuffer, + const std::shared_ptr &vkFence) +{ + VulkanCommandBufferList commandBufferList; + commandBufferList.add(commandBuffer); + + VkSubmitInfo vkSubmitInfo = {}; + vkSubmitInfo.sType = VK_STRUCTURE_TYPE_SUBMIT_INFO; + vkSubmitInfo.pNext = NULL; + vkSubmitInfo.waitSemaphoreCount = (uint32_t)0; + vkSubmitInfo.commandBufferCount = (uint32_t)commandBufferList.size(); + vkSubmitInfo.pCommandBuffers = commandBufferList(); + + vkQueueSubmit(m_vkQueue, 1, &vkSubmitInfo, vkFence->fence); +} + void VulkanQueue::submit(const VulkanSemaphoreList &waitSemaphoreList, const VulkanCommandBufferList &commandBufferList, const VulkanSemaphoreList &signalSemaphoreList) diff --git a/test_conformance/common/vulkan_wrapper/vulkan_wrapper.hpp b/test_conformance/common/vulkan_wrapper/vulkan_wrapper.hpp index 37925ee4a..af4782191 100644 --- a/test_conformance/common/vulkan_wrapper/vulkan_wrapper.hpp +++ b/test_conformance/common/vulkan_wrapper/vulkan_wrapper.hpp @@ -21,6 +21,7 @@ #include "vulkan_wrapper_types.hpp" #include "vulkan_list_map.hpp" #include "vulkan_api_list.hpp" +#include class VulkanInstance { friend const VulkanInstance &getVulkanInstance(); @@ -145,6 +146,20 @@ class VulkanDevice { operator VkDevice() const; }; +class VulkanFence { + friend class VulkanQueue; + +protected: + VkFence fence; + VkDevice device; + +public: + VulkanFence(const VulkanDevice &device); + virtual ~VulkanFence(); + void reset(); + void wait(); +}; + class VulkanQueue { friend class VulkanDevice; @@ -157,6 +172,8 @@ class VulkanQueue { public: const VulkanQueueFamily &getQueueFamily(); + void submit(const VulkanCommandBuffer &commandBuffer, + const std::shared_ptr &fence); void submit(const VulkanSemaphoreList &waitSemaphoreList, const VulkanCommandBufferList &commandBufferList, const VulkanSemaphoreList &signalSemaphoreList); @@ -569,7 +586,6 @@ class VulkanSemaphore { operator VkSemaphore() const; }; - #define VK_FUNC_DECL(name) extern "C" PFN_##name _##name; VK_FUNC_LIST #if defined(_WIN32) || defined(_WIN64) diff --git a/test_conformance/vulkan/main.cpp b/test_conformance/vulkan/main.cpp index 2eeb0c361..c37d61637 100644 --- a/test_conformance/vulkan/main.cpp +++ b/test_conformance/vulkan/main.cpp @@ -52,7 +52,8 @@ static void params_reset() } extern int test_buffer_common(cl_device_id device_, cl_context context_, - cl_command_queue queue_, int numElements_); + cl_command_queue queue_, int numElements_, + float use_fence); extern int test_image_common(cl_device_id device_, cl_context context_, cl_command_queue queue_, int numElements_); @@ -61,7 +62,7 @@ int test_buffer_single_queue(cl_device_id device_, cl_context context_, { params_reset(); log_info("RUNNING TEST WITH ONE QUEUE...... \n\n"); - return test_buffer_common(device_, context_, queue_, numElements_); + 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_) @@ -69,7 +70,7 @@ int test_buffer_multiple_queue(cl_device_id device_, cl_context context_, params_reset(); numCQ = 2; log_info("RUNNING TEST WITH TWO QUEUE...... \n\n"); - return test_buffer_common(device_, context_, queue_, numElements_); + 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_) @@ -78,7 +79,7 @@ int test_buffer_multiImport_sameCtx(cl_device_id device_, cl_context context_, multiImport = true; log_info("RUNNING TEST WITH MULTIPLE DEVICE MEMORY IMPORT " "IN SAME CONTEXT...... \n\n"); - return test_buffer_common(device_, context_, queue_, numElements_); + 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_) @@ -88,7 +89,45 @@ int test_buffer_multiImport_diffCtx(cl_device_id device_, cl_context context_, multiCtx = true; log_info("RUNNING TEST WITH MULTIPLE DEVICE MEMORY IMPORT " "IN DIFFERENT CONTEXT...... \n\n"); - return test_buffer_common(device_, context_, queue_, numElements_); + 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_) @@ -110,6 +149,10 @@ test_definition test_list[] = { ADD_TEST(buffer_single_queue), ADD_TEST(buffer_multiple_queue), ADD_TEST(buffer_multiImport_sameCtx), ADD_TEST(buffer_multiImport_diffCtx), + ADD_TEST(buffer_single_queue_fence), + ADD_TEST(buffer_multiple_queue_fence), + ADD_TEST(buffer_multiImport_sameCtx_fence), + ADD_TEST(buffer_multiImport_diffCtx_fence), ADD_TEST(image_single_queue), ADD_TEST(image_multiple_queue), ADD_TEST(consistency_external_buffer), diff --git a/test_conformance/vulkan/test_vulkan_interop_buffer.cpp b/test_conformance/vulkan/test_vulkan_interop_buffer.cpp index 9b0bc9de7..5390ef690 100644 --- a/test_conformance/vulkan/test_vulkan_interop_buffer.cpp +++ b/test_conformance/vulkan/test_vulkan_interop_buffer.cpp @@ -21,6 +21,7 @@ #include #include #include +#include #include #include "harness/errorHelpers.h" @@ -82,7 +83,8 @@ __kernel void checkKernel(__global unsigned char *ptr, int size, int expVal, __g int run_test_with_two_queue(cl_context &context, cl_command_queue &cmd_queue1, cl_command_queue &cmd_queue2, cl_kernel *kernel, cl_kernel &verify_kernel, VulkanDevice &vkDevice, - uint32_t numBuffers, uint32_t bufferSize) + uint32_t numBuffers, uint32_t bufferSize, + bool use_fence) { int err = CL_SUCCESS; size_t global_work_size[1]; @@ -117,6 +119,7 @@ int run_test_with_two_queue(cl_context &context, cl_command_queue &cmd_queue1, getSupportedVulkanExternalSemaphoreHandleTypeList()[0]; VulkanSemaphore vkVk2CLSemaphore(vkDevice, vkExternalSemaphoreHandleType); VulkanSemaphore vkCl2VkSemaphore(vkDevice, vkExternalSemaphoreHandleType); + std::shared_ptr fence = nullptr; VulkanQueue &vkQueue = vkDevice.getQueue(); @@ -136,10 +139,17 @@ int run_test_with_two_queue(cl_context &context, cl_command_queue &cmd_queue1, VulkanDescriptorSet vkDescriptorSet(vkDevice, vkDescriptorPool, vkDescriptorSetLayout); - clVk2CLExternalSemaphore = new clExternalSemaphore( - vkVk2CLSemaphore, context, vkExternalSemaphoreHandleType, deviceId); - clCl2VkExternalSemaphore = new clExternalSemaphore( - vkCl2VkSemaphore, context, vkExternalSemaphoreHandleType, deviceId); + if (use_fence) + { + fence = std::make_shared(vkDevice); + } + else + { + clVk2CLExternalSemaphore = new clExternalSemaphore( + vkVk2CLSemaphore, context, vkExternalSemaphoreHandleType, deviceId); + clCl2VkExternalSemaphore = new clExternalSemaphore( + vkCl2VkSemaphore, context, vkExternalSemaphoreHandleType, deviceId); + } const uint32_t maxIter = innerIterations; VulkanCommandPool vkCommandPool(vkDevice); @@ -227,16 +237,27 @@ int run_test_with_two_queue(cl_context &context, cl_command_queue &cmd_queue1, for (uint32_t iter = 0; iter < maxIter; iter++) { - if (iter == 0) + if (use_fence) { - vkQueue.submit(vkCommandBuffer, vkVk2CLSemaphore); + fence->reset(); + vkQueue.submit(vkCommandBuffer, fence); + fence->wait(); } else { - vkQueue.submit(vkCl2VkSemaphore, vkCommandBuffer, - vkVk2CLSemaphore); + if (iter == 0) + { + vkQueue.submit(vkCommandBuffer, vkVk2CLSemaphore); + } + else + { + vkQueue.submit(vkCl2VkSemaphore, vkCommandBuffer, + vkVk2CLSemaphore); + } + + clVk2CLExternalSemaphore->wait(cmd_queue1); } - clVk2CLExternalSemaphore->wait(cmd_queue1); + err = clSetKernelArg(update_buffer_kernel, 0, sizeof(uint32_t), (void *)&bufferSize); @@ -286,7 +307,14 @@ int run_test_with_two_queue(cl_context &context, cl_command_queue &cmd_queue1, goto CLEANUP; } - if (iter != (maxIter - 1)) + if (use_fence) + { + clFlush(cmd_queue1); + clFlush(cmd_queue2); + clFinish(cmd_queue1); + clFinish(cmd_queue2); + } + else if (!use_fence && iter != (maxIter - 1)) { clCl2VkExternalSemaphore->signal(cmd_queue2); } @@ -387,8 +415,11 @@ int run_test_with_two_queue(cl_context &context, cl_command_queue &cmd_queue1, } if (program) clReleaseProgram(program); if (kernel_cq) clReleaseKernel(kernel_cq); - if (clVk2CLExternalSemaphore) delete clVk2CLExternalSemaphore; - if (clCl2VkExternalSemaphore) delete clCl2VkExternalSemaphore; + if (!use_fence) + { + if (clVk2CLExternalSemaphore) delete clVk2CLExternalSemaphore; + if (clCl2VkExternalSemaphore) delete clCl2VkExternalSemaphore; + } if (error_2) free(error_2); if (error_1) clReleaseMemObject(error_1); @@ -398,7 +429,7 @@ int run_test_with_two_queue(cl_context &context, cl_command_queue &cmd_queue1, int run_test_with_one_queue(cl_context &context, cl_command_queue &cmd_queue1, cl_kernel *kernel, cl_kernel &verify_kernel, VulkanDevice &vkDevice, uint32_t numBuffers, - uint32_t bufferSize) + uint32_t bufferSize, bool use_fence) { log_info("RUNNING TEST WITH ONE QUEUE...... \n\n"); size_t global_work_size[1]; @@ -416,6 +447,7 @@ int run_test_with_one_queue(cl_context &context, cl_command_queue &cmd_queue1, getSupportedVulkanExternalSemaphoreHandleTypeList()[0]; VulkanSemaphore vkVk2CLSemaphore(vkDevice, vkExternalSemaphoreHandleType); VulkanSemaphore vkCl2VkSemaphore(vkDevice, vkExternalSemaphoreHandleType); + std::shared_ptr fence = nullptr; VulkanQueue &vkQueue = vkDevice.getQueue(); @@ -434,10 +466,18 @@ int run_test_with_one_queue(cl_context &context, cl_command_queue &cmd_queue1, VulkanDescriptorSet vkDescriptorSet(vkDevice, vkDescriptorPool, vkDescriptorSetLayout); - clVk2CLExternalSemaphore = new clExternalSemaphore( - vkVk2CLSemaphore, context, vkExternalSemaphoreHandleType, deviceId); - clCl2VkExternalSemaphore = new clExternalSemaphore( - vkCl2VkSemaphore, context, vkExternalSemaphoreHandleType, deviceId); + if (use_fence) + { + fence = std::make_shared(vkDevice); + } + else + { + clVk2CLExternalSemaphore = new clExternalSemaphore( + vkVk2CLSemaphore, context, vkExternalSemaphoreHandleType, deviceId); + clCl2VkExternalSemaphore = new clExternalSemaphore( + vkCl2VkSemaphore, context, vkExternalSemaphoreHandleType, deviceId); + } + const uint32_t maxIter = innerIterations; VulkanCommandPool vkCommandPool(vkDevice); VulkanCommandBuffer vkCommandBuffer(vkDevice, vkCommandPool); @@ -526,16 +566,26 @@ int run_test_with_one_queue(cl_context &context, cl_command_queue &cmd_queue1, for (uint32_t iter = 0; iter < maxIter; iter++) { - if (iter == 0) + if (use_fence) { - vkQueue.submit(vkCommandBuffer, vkVk2CLSemaphore); + fence->reset(); + vkQueue.submit(vkCommandBuffer, fence); + fence->wait(); } else { - vkQueue.submit(vkCl2VkSemaphore, vkCommandBuffer, - vkVk2CLSemaphore); + if (iter == 0) + { + vkQueue.submit(vkCommandBuffer, vkVk2CLSemaphore); + } + else + { + vkQueue.submit(vkCl2VkSemaphore, vkCommandBuffer, + vkVk2CLSemaphore); + } + + clVk2CLExternalSemaphore->wait(cmd_queue1); } - clVk2CLExternalSemaphore->wait(cmd_queue1); err = clSetKernelArg(update_buffer_kernel, 0, sizeof(uint32_t), (void *)&bufferSize); @@ -562,7 +612,12 @@ int run_test_with_one_queue(cl_context &context, cl_command_queue &cmd_queue1, " error\n"); goto CLEANUP; } - if (iter != (maxIter - 1)) + if (use_fence) + { + clFlush(cmd_queue1); + clFinish(cmd_queue1); + } + else if (!use_fence && (iter != (maxIter - 1))) { clCl2VkExternalSemaphore->signal(cmd_queue1); } @@ -656,8 +711,13 @@ int run_test_with_one_queue(cl_context &context, cl_command_queue &cmd_queue1, delete externalMemory[i]; } } - if (clVk2CLExternalSemaphore) delete clVk2CLExternalSemaphore; - if (clCl2VkExternalSemaphore) delete clCl2VkExternalSemaphore; + + if (!use_fence) + { + if (clVk2CLExternalSemaphore) delete clVk2CLExternalSemaphore; + if (clCl2VkExternalSemaphore) delete clCl2VkExternalSemaphore; + } + if (error_2) free(error_2); if (error_1) clReleaseMemObject(error_1); return err; @@ -666,7 +726,7 @@ int run_test_with_one_queue(cl_context &context, cl_command_queue &cmd_queue1, int run_test_with_multi_import_same_ctx( cl_context &context, cl_command_queue &cmd_queue1, cl_kernel *kernel, cl_kernel &verify_kernel, VulkanDevice &vkDevice, uint32_t numBuffers, - uint32_t bufferSize, uint32_t bufferSizeForOffset) + uint32_t bufferSize, uint32_t bufferSizeForOffset, float use_fence) { size_t global_work_size[1]; uint8_t *error_2; @@ -687,6 +747,7 @@ int run_test_with_multi_import_same_ctx( getSupportedVulkanExternalSemaphoreHandleTypeList()[0]; VulkanSemaphore vkVk2CLSemaphore(vkDevice, vkExternalSemaphoreHandleType); VulkanSemaphore vkCl2VkSemaphore(vkDevice, vkExternalSemaphoreHandleType); + std::shared_ptr fence = nullptr; VulkanQueue &vkQueue = vkDevice.getQueue(); @@ -706,10 +767,18 @@ int run_test_with_multi_import_same_ctx( VulkanDescriptorSet vkDescriptorSet(vkDevice, vkDescriptorPool, vkDescriptorSetLayout); - clVk2CLExternalSemaphore = new clExternalSemaphore( - vkVk2CLSemaphore, context, vkExternalSemaphoreHandleType, deviceId); - clCl2VkExternalSemaphore = new clExternalSemaphore( - vkCl2VkSemaphore, context, vkExternalSemaphoreHandleType, deviceId); + if (use_fence) + { + fence = std::make_shared(vkDevice); + } + else + { + clVk2CLExternalSemaphore = new clExternalSemaphore( + vkVk2CLSemaphore, context, vkExternalSemaphoreHandleType, deviceId); + clCl2VkExternalSemaphore = new clExternalSemaphore( + vkCl2VkSemaphore, context, vkExternalSemaphoreHandleType, deviceId); + } + const uint32_t maxIter = innerIterations; VulkanCommandPool vkCommandPool(vkDevice); VulkanCommandBuffer vkCommandBuffer(vkDevice, vkCommandPool); @@ -832,16 +901,34 @@ int run_test_with_multi_import_same_ctx( for (uint32_t iter = 0; iter < maxIter; iter++) { - if (iter == 0) + if (use_fence) { - vkQueue.submit(vkCommandBuffer, vkVk2CLSemaphore); + fence->reset(); + vkQueue.submit(vkCommandBuffer, fence); + fence->wait(); } else { - vkQueue.submit(vkCl2VkSemaphore, vkCommandBuffer, - vkVk2CLSemaphore); + if (iter == 0) + { + vkQueue.submit(vkCommandBuffer, vkVk2CLSemaphore); + } + else + { + vkQueue.submit(vkCl2VkSemaphore, vkCommandBuffer, + vkVk2CLSemaphore); + } } - clVk2CLExternalSemaphore->wait(cmd_queue1); + + if (use_fence) + { + fence->wait(); + } + else + { + clVk2CLExternalSemaphore->wait(cmd_queue1); + } + for (uint8_t launchIter = 0; launchIter < numImports; launchIter++) { @@ -874,7 +961,11 @@ int run_test_with_multi_import_same_ctx( goto CLEANUP; } } - if (iter != (maxIter - 1)) + if (use_fence) + { + clFinish(cmd_queue1); + } + else if (!use_fence && iter != (maxIter - 1)) { clCl2VkExternalSemaphore->signal(cmd_queue1); } @@ -987,8 +1078,13 @@ int run_test_with_multi_import_same_ctx( } } } - if (clVk2CLExternalSemaphore) delete clVk2CLExternalSemaphore; - if (clCl2VkExternalSemaphore) delete clCl2VkExternalSemaphore; + + if (!use_fence) + { + if (clVk2CLExternalSemaphore) delete clVk2CLExternalSemaphore; + if (clCl2VkExternalSemaphore) delete clCl2VkExternalSemaphore; + } + if (error_2) free(error_2); if (error_1) clReleaseMemObject(error_1); return err; @@ -998,7 +1094,8 @@ int run_test_with_multi_import_diff_ctx( cl_context &context, cl_context &context2, cl_command_queue &cmd_queue1, cl_command_queue &cmd_queue2, cl_kernel *kernel1, cl_kernel *kernel2, cl_kernel &verify_kernel, cl_kernel verify_kernel2, VulkanDevice &vkDevice, - uint32_t numBuffers, uint32_t bufferSize, uint32_t bufferSizeForOffset) + uint32_t numBuffers, uint32_t bufferSize, uint32_t bufferSizeForOffset, + float use_fence) { size_t global_work_size[1]; uint8_t *error_3; @@ -1023,6 +1120,7 @@ int run_test_with_multi_import_diff_ctx( getSupportedVulkanExternalSemaphoreHandleTypeList()[0]; VulkanSemaphore vkVk2CLSemaphore(vkDevice, vkExternalSemaphoreHandleType); VulkanSemaphore vkCl2VkSemaphore(vkDevice, vkExternalSemaphoreHandleType); + std::shared_ptr fence = nullptr; VulkanQueue &vkQueue = vkDevice.getQueue(); @@ -1042,15 +1140,24 @@ int run_test_with_multi_import_diff_ctx( VulkanDescriptorSet vkDescriptorSet(vkDevice, vkDescriptorPool, vkDescriptorSetLayout); - clVk2CLExternalSemaphore = new clExternalSemaphore( - vkVk2CLSemaphore, context, vkExternalSemaphoreHandleType, deviceId); - clCl2VkExternalSemaphore = new clExternalSemaphore( - vkCl2VkSemaphore, context, vkExternalSemaphoreHandleType, deviceId); - - clVk2CLExternalSemaphore2 = new clExternalSemaphore( - vkVk2CLSemaphore, context2, vkExternalSemaphoreHandleType, deviceId); - clCl2VkExternalSemaphore2 = new clExternalSemaphore( - vkCl2VkSemaphore, context2, vkExternalSemaphoreHandleType, deviceId); + if (use_fence) + { + fence = std::make_shared(vkDevice); + } + else + { + clVk2CLExternalSemaphore = new clExternalSemaphore( + vkVk2CLSemaphore, context, vkExternalSemaphoreHandleType, deviceId); + clCl2VkExternalSemaphore = new clExternalSemaphore( + vkCl2VkSemaphore, context, vkExternalSemaphoreHandleType, deviceId); + + clVk2CLExternalSemaphore2 = + new clExternalSemaphore(vkVk2CLSemaphore, context2, + vkExternalSemaphoreHandleType, deviceId); + clCl2VkExternalSemaphore2 = + new clExternalSemaphore(vkCl2VkSemaphore, context2, + vkExternalSemaphoreHandleType, deviceId); + } const uint32_t maxIter = innerIterations; VulkanCommandPool vkCommandPool(vkDevice); @@ -1192,16 +1299,33 @@ int run_test_with_multi_import_diff_ctx( for (uint32_t iter = 0; iter < maxIter; iter++) { - if (iter == 0) + if (use_fence) { - vkQueue.submit(vkCommandBuffer, vkVk2CLSemaphore); + fence->reset(); + vkQueue.submit(vkCommandBuffer, fence); + fence->wait(); } else { - vkQueue.submit(vkCl2VkSemaphore, vkCommandBuffer, - vkVk2CLSemaphore); + if (iter == 0) + { + vkQueue.submit(vkCommandBuffer, vkVk2CLSemaphore); + } + else + { + vkQueue.submit(vkCl2VkSemaphore, vkCommandBuffer, + vkVk2CLSemaphore); + } + } + + if (use_fence) + { + fence->wait(); + } + else + { + clVk2CLExternalSemaphore->wait(cmd_queue1); } - clVk2CLExternalSemaphore->wait(cmd_queue1); for (uint8_t launchIter = 0; launchIter < numImports; launchIter++) @@ -1235,7 +1359,11 @@ int run_test_with_multi_import_diff_ctx( goto CLEANUP; } } - if (iter != (maxIter - 1)) + if (use_fence) + { + clFinish(cmd_queue1); + } + else if (!use_fence && iter != (maxIter - 1)) { clCl2VkExternalSemaphore->signal(cmd_queue1); } @@ -1243,16 +1371,33 @@ int run_test_with_multi_import_diff_ctx( clFinish(cmd_queue1); for (uint32_t iter = 0; iter < maxIter; iter++) { - if (iter == 0) + if (use_fence) { - vkQueue.submit(vkCommandBuffer, vkVk2CLSemaphore); + fence->reset(); + vkQueue.submit(vkCommandBuffer, fence); + fence->wait(); } else { - vkQueue.submit(vkCl2VkSemaphore, vkCommandBuffer, - vkVk2CLSemaphore); + if (iter == 0) + { + vkQueue.submit(vkCommandBuffer, vkVk2CLSemaphore); + } + else + { + vkQueue.submit(vkCl2VkSemaphore, vkCommandBuffer, + vkVk2CLSemaphore); + } + } + + if (use_fence) + { + fence->wait(); + } + else + { + clVk2CLExternalSemaphore2->wait(cmd_queue2); } - clVk2CLExternalSemaphore2->wait(cmd_queue2); for (uint8_t launchIter = 0; launchIter < numImports; launchIter++) @@ -1286,7 +1431,11 @@ int run_test_with_multi_import_diff_ctx( goto CLEANUP; } } - if (iter != (maxIter - 1)) + if (use_fence) + { + clFinish(cmd_queue2); + } + else if (!use_fence && iter != (maxIter - 1)) { clCl2VkExternalSemaphore2->signal(cmd_queue2); } @@ -1474,10 +1623,15 @@ int run_test_with_multi_import_diff_ctx( } } } - if (clVk2CLExternalSemaphore) delete clVk2CLExternalSemaphore; - if (clCl2VkExternalSemaphore) delete clCl2VkExternalSemaphore; - if (clVk2CLExternalSemaphore2) delete clVk2CLExternalSemaphore2; - if (clCl2VkExternalSemaphore2) delete clCl2VkExternalSemaphore2; + + if (!use_fence) + { + if (clVk2CLExternalSemaphore) delete clVk2CLExternalSemaphore; + if (clCl2VkExternalSemaphore) delete clCl2VkExternalSemaphore; + if (clVk2CLExternalSemaphore2) delete clVk2CLExternalSemaphore2; + if (clCl2VkExternalSemaphore2) delete clCl2VkExternalSemaphore2; + } + if (error_3) free(error_3); if (error_1) clReleaseMemObject(error_1); if (error_2) clReleaseMemObject(error_2); @@ -1485,7 +1639,8 @@ int run_test_with_multi_import_diff_ctx( } int test_buffer_common(cl_device_id device_, cl_context context_, - cl_command_queue queue_, int numElements_) + cl_command_queue queue_, int numElements_, + float use_fence) { int current_device = 0; @@ -1738,26 +1893,26 @@ int test_buffer_common(cl_device_id device_, cl_context context_, { errNum = run_test_with_multi_import_same_ctx( context, cmd_queue1, kernel, verify_kernel, vkDevice, - numBuffers, bufferSize, bufferSizeForOffset); + numBuffers, bufferSize, bufferSizeForOffset, use_fence); } else if (multiImport && multiCtx) { errNum = run_test_with_multi_import_diff_ctx( context, context2, cmd_queue1, cmd_queue3, kernel, kernel2, verify_kernel, verify_kernel2, vkDevice, numBuffers, - bufferSize, bufferSizeForOffset); + bufferSize, bufferSizeForOffset, use_fence); } else if (numCQ == 2) { errNum = run_test_with_two_queue( context, cmd_queue1, cmd_queue2, kernel, verify_kernel, - vkDevice, numBuffers + 1, bufferSize); + vkDevice, numBuffers + 1, bufferSize, use_fence); } else { - errNum = run_test_with_one_queue(context, cmd_queue1, kernel, - verify_kernel, vkDevice, - numBuffers, bufferSize); + errNum = run_test_with_one_queue( + context, cmd_queue1, kernel, verify_kernel, vkDevice, + numBuffers, bufferSize, use_fence); } if (errNum != CL_SUCCESS) {