From 714b54c6b77b3692a9ca06b2b173f42556447e28 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Pawe=C5=82=20Jastrz=C4=99bski?= Date: Tue, 21 Feb 2023 11:52:58 +0100 Subject: [PATCH 01/14] Add tests for external sharing not dependant on semaphores. MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Additional external sharing tests that use fences instead of semaphores. Signed-off-by: Paweł Jastrzębski --- test_conformance/vulkan/CMakeLists.txt | 1 + test_conformance/vulkan/main.cpp | 69 +- .../vulkan/test_vulkan_interop_buffer.cpp | 8 +- .../test_vulkan_interop_buffer_fence.cpp | 1798 +++++++++++++++++ .../vulkan_interop_common/vulkan_wrapper.cpp | 10 + .../vulkan_interop_common/vulkan_wrapper.hpp | 1 + 6 files changed, 1872 insertions(+), 15 deletions(-) create mode 100644 test_conformance/vulkan/test_vulkan_interop_buffer_fence.cpp diff --git a/test_conformance/vulkan/CMakeLists.txt b/test_conformance/vulkan/CMakeLists.txt index 4f43172af..004a2c0a8 100644 --- a/test_conformance/vulkan/CMakeLists.txt +++ b/test_conformance/vulkan/CMakeLists.txt @@ -23,6 +23,7 @@ include_directories (${CLConform_INCLUDE_DIR}) set (${MODULE_NAME}_SOURCES main.cpp test_vulkan_interop_buffer.cpp + test_vulkan_interop_buffer_fence.cpp test_vulkan_interop_image.cpp test_vulkan_api_consistency.cpp test_vulkan_platform_device_info.cpp diff --git a/test_conformance/vulkan/main.cpp b/test_conformance/vulkan/main.cpp index 2eeb0c361..f8dad951a 100644 --- a/test_conformance/vulkan/main.cpp +++ b/test_conformance/vulkan/main.cpp @@ -53,6 +53,8 @@ static void params_reset() extern int test_buffer_common(cl_device_id device_, cl_context context_, cl_command_queue queue_, int numElements_); +extern int test_buffer_common_fence(cl_device_id device_, cl_context context_, + cl_command_queue queue_, int numElements_); extern int test_image_common(cl_device_id device_, cl_context context_, cl_command_queue queue_, int numElements_); @@ -90,6 +92,44 @@ int test_buffer_multiImport_diffCtx(cl_device_id device_, cl_context context_, "IN DIFFERENT CONTEXT...... \n\n"); return test_buffer_common(device_, context_, queue_, numElements_); } +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_fence(device_, context_, queue_, numElements_); +} +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_fence(device_, context_, queue_, numElements_); +} +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_fence(device_, context_, queue_, numElements_); +} +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_fence(device_, context_, queue_, numElements_); +} int test_image_single_queue(cl_device_id device_, cl_context context_, cl_command_queue queue_, int numElements_) { @@ -106,17 +146,24 @@ int test_image_multiple_queue(cl_device_id device_, cl_context context_, 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), - ADD_TEST(buffer_multiImport_diffCtx), - ADD_TEST(image_single_queue), - ADD_TEST(image_multiple_queue), - ADD_TEST(consistency_external_buffer), - ADD_TEST(consistency_external_image), - ADD_TEST(consistency_external_semaphore), - ADD_TEST(platform_info), - ADD_TEST(device_info) }; +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), + ADD_TEST(consistency_external_image), + ADD_TEST(consistency_external_semaphore), + ADD_TEST(platform_info), + ADD_TEST(device_info) +}; const int test_num = ARRAY_SIZE(test_list); diff --git a/test_conformance/vulkan/test_vulkan_interop_buffer.cpp b/test_conformance/vulkan/test_vulkan_interop_buffer.cpp index 9b0bc9de7..7d3062b3f 100644 --- a/test_conformance/vulkan/test_vulkan_interop_buffer.cpp +++ b/test_conformance/vulkan/test_vulkan_interop_buffer.cpp @@ -39,7 +39,7 @@ struct Params }; } -const char *kernel_text_numbuffer_1 = " \ +static 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\ @@ -47,7 +47,7 @@ __kernel void clUpdateBuffer(int bufferSize, __global unsigned char *a) { \n\ } \n\ }"; -const char *kernel_text_numbuffer_2 = " \ +static 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\ @@ -56,7 +56,7 @@ __kernel void clUpdateBuffer(int bufferSize, __global unsigned char *a, __global } \n\ }"; -const char *kernel_text_numbuffer_4 = " \ +static 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\ @@ -68,7 +68,7 @@ __kernel void clUpdateBuffer(int bufferSize, __global unsigned char *a, __global }"; -const char *kernel_text_verify = " \ +static 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\ diff --git a/test_conformance/vulkan/test_vulkan_interop_buffer_fence.cpp b/test_conformance/vulkan/test_vulkan_interop_buffer_fence.cpp new file mode 100644 index 000000000..c841e612e --- /dev/null +++ b/test_conformance/vulkan/test_vulkan_interop_buffer_fence.cpp @@ -0,0 +1,1798 @@ +// +// Copyright (c) 2022 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 +#include +#include +#include +#include +#include +#include "harness/errorHelpers.h" + +#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 { +struct Params +{ + uint32_t numBuffers; + uint32_t bufferSize; + uint32_t interBufferOffset; +}; +} + +static 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\ +}"; + +static 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\ +}"; + +static 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\ +}"; + + +static 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\ +}"; + +int run_test_with_two_queue_fence(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) +{ + int err = CL_SUCCESS; + size_t global_work_size[1]; + uint8_t *error_2; + cl_mem error_1; + cl_kernel update_buffer_kernel; + cl_kernel kernel_cq; + const char *program_source_const = kernel_text_numbuffer_2; + size_t program_source_length = strlen(program_source_const); + cl_program program = clCreateProgramWithSource( + context, 1, &program_source_const, &program_source_length, &err); + err = clBuildProgram(program, 0, NULL, NULL, NULL, NULL); + if (err != CL_SUCCESS) + { + print_error(err, "Error: Failed to build program \n"); + return err; + } + // create the kernel + kernel_cq = clCreateKernel(program, "clUpdateBuffer", &err); + if (err != CL_SUCCESS) + { + print_error(err, "clCreateKernel failed \n"); + return err; + } + + const std::vector + vkExternalMemoryHandleTypeList = + getSupportedVulkanExternalMemoryHandleTypeList(); + + VulkanQueue &vkQueue = vkDevice.getQueue(); + + std::vector vkBufferShader = readFile("buffer.spv"); + + VulkanShaderModule vkBufferShaderModule(vkDevice, vkBufferShader); + VulkanDescriptorSetLayoutBindingList vkDescriptorSetLayoutBindingList( + MAX_BUFFERS + 1, VULKAN_DESCRIPTOR_TYPE_STORAGE_BUFFER); + VulkanDescriptorSetLayout vkDescriptorSetLayout( + vkDevice, vkDescriptorSetLayoutBindingList); + VulkanPipelineLayout vkPipelineLayout(vkDevice, vkDescriptorSetLayout); + VulkanComputePipeline vkComputePipeline(vkDevice, vkPipelineLayout, + vkBufferShaderModule); + + VulkanDescriptorPool vkDescriptorPool(vkDevice, + vkDescriptorSetLayoutBindingList); + VulkanDescriptorSet vkDescriptorSet(vkDevice, vkDescriptorPool, + vkDescriptorSetLayout); + + const uint32_t maxIter = innerIterations; + VulkanCommandPool vkCommandPool(vkDevice); + VulkanCommandBuffer vkCommandBuffer(vkDevice, vkCommandPool); + + VulkanBuffer vkParamsBuffer(vkDevice, sizeof(Params)); + VulkanDeviceMemory vkParamsDeviceMemory( + vkDevice, vkParamsBuffer.getSize(), + getVulkanMemoryType(vkDevice, + VULKAN_MEMORY_TYPE_PROPERTY_HOST_VISIBLE_COHERENT)); + vkParamsDeviceMemory.bindBuffer(vkParamsBuffer); + std::vector vkBufferListDeviceMemory; + std::vector externalMemory; + + VkFence fence; + VkFenceCreateInfo fenceInfo{}; + fenceInfo.sType = VK_STRUCTURE_TYPE_FENCE_CREATE_INFO; + fenceInfo.pNext = nullptr; + fenceInfo.flags = 0; + + VkResult vkStatus = vkCreateFence(vkDevice, &fenceInfo, nullptr, &fence); + + if (vkStatus != VK_SUCCESS) + { + print_error(vkStatus, "Error: Failed create fence.\n"); + goto CLEANUP; + } + + for (size_t emhtIdx = 0; emhtIdx < vkExternalMemoryHandleTypeList.size(); + emhtIdx++) + { + VulkanExternalMemoryHandleType vkExternalMemoryHandleType = + vkExternalMemoryHandleTypeList[emhtIdx]; + log_info("External memory handle type: %d\n", + vkExternalMemoryHandleType); + + VulkanBuffer vkDummyBuffer(vkDevice, 4 * 1024, + vkExternalMemoryHandleType); + const VulkanMemoryTypeList &memoryTypeList = + vkDummyBuffer.getMemoryTypeList(); + + for (size_t mtIdx = 0; mtIdx < memoryTypeList.size(); mtIdx++) + { + const VulkanMemoryType &memoryType = memoryTypeList[mtIdx]; + + log_info("Memory type index: %d\n", (uint32_t)memoryType); + log_info("Memory type property: %d\n", + memoryType.getMemoryTypeProperty()); + + VulkanBufferList vkBufferList(numBuffers, vkDevice, bufferSize, + vkExternalMemoryHandleType); + + for (size_t bIdx = 0; bIdx < numBuffers; bIdx++) + { + vkBufferListDeviceMemory.push_back( + new VulkanDeviceMemory(vkDevice, bufferSize, memoryType, + vkExternalMemoryHandleType)); + externalMemory.push_back(new clExternalMemory( + vkBufferListDeviceMemory[bIdx], vkExternalMemoryHandleType, + 0, bufferSize, context, deviceId)); + } + cl_mem buffers[MAX_BUFFERS]; + 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); + buffers[bIdx] = externalMemory[bIdx]->getExternalMemoryBuffer(); + vkDescriptorSet.update((uint32_t)bIdx + 1, vkBufferList[bIdx]); + } + vkCommandBuffer.begin(); + vkCommandBuffer.bindPipeline(vkComputePipeline); + vkCommandBuffer.bindDescriptorSets( + vkComputePipeline, vkPipelineLayout, vkDescriptorSet); + vkCommandBuffer.dispatch(512, 1, 1); + vkCommandBuffer.end(); + + if (vkBufferList.size() == 2) + { + update_buffer_kernel = kernel[0]; + } + else if (vkBufferList.size() == 3) + { + update_buffer_kernel = kernel[1]; + } + else if (vkBufferList.size() == 5) + { + update_buffer_kernel = 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++) + { + + if (iter == 0) + { + vkQueue.submit(vkCommandBuffer, fence); + } + else + { + vkQueue.submit(vkCommandBuffer, fence); + } + vkWaitForFences(vkDevice, 1, &fence, VK_TRUE, UINT64_MAX); + + err = clSetKernelArg(update_buffer_kernel, 0, sizeof(uint32_t), + (void *)&bufferSize); + err |= clSetKernelArg(kernel_cq, 0, sizeof(uint32_t), + (void *)&bufferSize); + err |= clSetKernelArg(kernel_cq, 1, sizeof(cl_mem), + (void *)&(buffers[0])); + + for (int i = 0; i < vkBufferList.size() - 1; i++) + { + err |= + clSetKernelArg(update_buffer_kernel, i + 1, + sizeof(cl_mem), (void *)&(buffers[i])); + } + + err |= + clSetKernelArg(kernel_cq, 2, sizeof(cl_mem), + (void *)&(buffers[vkBufferList.size() - 1])); + + if (err != CL_SUCCESS) + { + print_error(err, + "Error: Failed to set arg values for kernel\n"); + goto CLEANUP; + } + cl_event first_launch; + + err = clEnqueueNDRangeKernel(cmd_queue1, update_buffer_kernel, + 1, NULL, global_work_size, NULL, 0, + NULL, &first_launch); + if (err != CL_SUCCESS) + { + print_error(err, + "Error: Failed to launch update_buffer_kernel," + "error\n"); + goto CLEANUP; + } + + err = clEnqueueNDRangeKernel(cmd_queue2, kernel_cq, 1, NULL, + global_work_size, NULL, 1, + &first_launch, NULL); + if (err != CL_SUCCESS) + { + print_error(err, + "Error: Failed to launch update_buffer_kernel," + "error\n"); + goto CLEANUP; + } + + if (iter != (maxIter - 1)) + { + vkWaitForFences(vkDevice, 1, &fence, VK_TRUE, UINT64_MAX); + } + } + error_2 = (uint8_t *)malloc(sizeof(uint8_t)); + if (NULL == error_2) + { + log_error("Not able to allocate memory\n"); + goto CLEANUP; + } + clFinish(cmd_queue2); + error_1 = clCreateBuffer(context, CL_MEM_WRITE_ONLY, + sizeof(uint8_t), NULL, &err); + if (CL_SUCCESS != err) + { + print_error(err, "Error: clCreateBuffer \n"); + goto CLEANUP; + } + uint8_t val = 0; + err = clEnqueueWriteBuffer(cmd_queue1, error_1, CL_TRUE, 0, + sizeof(uint8_t), &val, 0, NULL, NULL); + if (err != CL_SUCCESS) + { + print_error(err, "Error: Failed read output, error\n"); + goto CLEANUP; + } + + int calc_max_iter; + for (int i = 0; i < vkBufferList.size(); i++) + { + if (i == 0) + calc_max_iter = (maxIter * 3); + else + calc_max_iter = (maxIter * 2); + err = clSetKernelArg(verify_kernel, 0, sizeof(cl_mem), + (void *)&(buffers[i])); + 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); + if (err != CL_SUCCESS) + { + print_error(err, + "Error: Failed to set arg values for " + "verify_kernel \n"); + goto CLEANUP; + } + err = clEnqueueNDRangeKernel(cmd_queue1, verify_kernel, 1, NULL, + global_work_size, NULL, 0, NULL, + NULL); + + if (err != CL_SUCCESS) + { + print_error(err, + "Error: Failed to launch verify_kernel," + "error \n"); + goto CLEANUP; + } + err = clEnqueueReadBuffer(cmd_queue1, error_1, CL_TRUE, 0, + sizeof(uint8_t), error_2, 0, NULL, + NULL); + if (err != CL_SUCCESS) + { + print_error(err, "Error: Failed read output, error \n "); + goto CLEANUP; + } + if (*error_2 == 1) + { + log_error("&&&& vulkan_opencl_buffer test FAILED\n"); + goto CLEANUP; + } + } + for (size_t i = 0; i < vkBufferList.size(); i++) + { + delete vkBufferListDeviceMemory[i]; + delete externalMemory[i]; + } + vkBufferListDeviceMemory.erase(vkBufferListDeviceMemory.begin(), + vkBufferListDeviceMemory.begin() + + numBuffers); + externalMemory.erase(externalMemory.begin(), + externalMemory.begin() + numBuffers); + } + } +CLEANUP: + for (size_t i = 0; i < vkBufferListDeviceMemory.size(); i++) + { + if (vkBufferListDeviceMemory[i]) + { + delete vkBufferListDeviceMemory[i]; + } + if (externalMemory[i]) + { + delete externalMemory[i]; + } + } + if (program) clReleaseProgram(program); + if (kernel_cq) clReleaseKernel(kernel_cq); + if (error_2) free(error_2); + if (error_1) clReleaseMemObject(error_1); + + vkDestroyFence(vkDevice, fence, nullptr); + + return err; +} + +int run_test_with_one_queue_fence(cl_context &context, + cl_command_queue &cmd_queue1, + cl_kernel *kernel, cl_kernel &verify_kernel, + VulkanDevice &vkDevice, uint32_t numBuffers, + uint32_t bufferSize) +{ + log_info("RUNNING TEST WITH ONE QUEUE...... \n\n"); + size_t global_work_size[1]; + uint8_t *error_2; + cl_mem error_1; + cl_kernel update_buffer_kernel; + int err = CL_SUCCESS; + + const std::vector + vkExternalMemoryHandleTypeList = + getSupportedVulkanExternalMemoryHandleTypeList(); + + VulkanQueue &vkQueue = vkDevice.getQueue(); + + std::vector vkBufferShader = readFile("buffer.spv"); + VulkanShaderModule vkBufferShaderModule(vkDevice, vkBufferShader); + VulkanDescriptorSetLayoutBindingList vkDescriptorSetLayoutBindingList( + MAX_BUFFERS + 1, VULKAN_DESCRIPTOR_TYPE_STORAGE_BUFFER); + VulkanDescriptorSetLayout vkDescriptorSetLayout( + vkDevice, vkDescriptorSetLayoutBindingList); + VulkanPipelineLayout vkPipelineLayout(vkDevice, vkDescriptorSetLayout); + VulkanComputePipeline vkComputePipeline(vkDevice, vkPipelineLayout, + vkBufferShaderModule); + + VulkanDescriptorPool vkDescriptorPool(vkDevice, + vkDescriptorSetLayoutBindingList); + VulkanDescriptorSet vkDescriptorSet(vkDevice, vkDescriptorPool, + vkDescriptorSetLayout); + + const uint32_t maxIter = innerIterations; + VulkanCommandPool vkCommandPool(vkDevice); + VulkanCommandBuffer vkCommandBuffer(vkDevice, vkCommandPool); + + VulkanBuffer vkParamsBuffer(vkDevice, sizeof(Params)); + VulkanDeviceMemory vkParamsDeviceMemory( + vkDevice, vkParamsBuffer.getSize(), + getVulkanMemoryType(vkDevice, + VULKAN_MEMORY_TYPE_PROPERTY_HOST_VISIBLE_COHERENT)); + vkParamsDeviceMemory.bindBuffer(vkParamsBuffer); + std::vector vkBufferListDeviceMemory; + std::vector externalMemory; + + VkFence fence; + VkFenceCreateInfo fenceInfo{}; + fenceInfo.sType = VK_STRUCTURE_TYPE_FENCE_CREATE_INFO; + fenceInfo.pNext = nullptr; + fenceInfo.flags = 0; + + VkResult vkStatus = vkCreateFence(vkDevice, &fenceInfo, nullptr, &fence); + + if (vkStatus != VK_SUCCESS) + { + print_error(vkStatus, "Error: Failed create fence.\n"); + goto CLEANUP; + } + + for (size_t emhtIdx = 0; emhtIdx < vkExternalMemoryHandleTypeList.size(); + emhtIdx++) + { + VulkanExternalMemoryHandleType vkExternalMemoryHandleType = + vkExternalMemoryHandleTypeList[emhtIdx]; + log_info("External memory handle type: %d\n", + vkExternalMemoryHandleType); + + VulkanBuffer vkDummyBuffer(vkDevice, 4 * 1024, + vkExternalMemoryHandleType); + const VulkanMemoryTypeList &memoryTypeList = + vkDummyBuffer.getMemoryTypeList(); + + for (size_t mtIdx = 0; mtIdx < memoryTypeList.size(); mtIdx++) + { + const VulkanMemoryType &memoryType = memoryTypeList[mtIdx]; + + log_info("Memory type index: %d\n", (uint32_t)memoryType); + log_info("Memory type property: %d\n", + memoryType.getMemoryTypeProperty()); + + VulkanBufferList vkBufferList(numBuffers, vkDevice, bufferSize, + vkExternalMemoryHandleType); + + for (size_t bIdx = 0; bIdx < numBuffers; bIdx++) + { + vkBufferListDeviceMemory.push_back( + new VulkanDeviceMemory(vkDevice, bufferSize, memoryType, + vkExternalMemoryHandleType)); + externalMemory.push_back(new clExternalMemory( + vkBufferListDeviceMemory[bIdx], vkExternalMemoryHandleType, + 0, bufferSize, context, deviceId)); + } + cl_mem buffers[4]; + 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); + buffers[bIdx] = externalMemory[bIdx]->getExternalMemoryBuffer(); + vkDescriptorSet.update((uint32_t)bIdx + 1, vkBufferList[bIdx]); + } + vkCommandBuffer.begin(); + vkCommandBuffer.bindPipeline(vkComputePipeline); + vkCommandBuffer.bindDescriptorSets( + vkComputePipeline, vkPipelineLayout, vkDescriptorSet); + vkCommandBuffer.dispatch(512, 1, 1); + vkCommandBuffer.end(); + + if (vkBufferList.size() == 1) + { + update_buffer_kernel = kernel[0]; + } + else if (vkBufferList.size() == 2) + { + update_buffer_kernel = kernel[1]; + } + else if (vkBufferList.size() == 4) + { + update_buffer_kernel = 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++) + { + if (iter == 0) + { + vkQueue.submit(vkCommandBuffer, fence); + } + else + { + vkQueue.submit(vkCommandBuffer, fence); + } + vkWaitForFences(vkDevice, 1, &fence, VK_TRUE, UINT64_MAX); + + err = clSetKernelArg(update_buffer_kernel, 0, sizeof(uint32_t), + (void *)&bufferSize); + for (int i = 0; i < vkBufferList.size(); i++) + { + err |= + clSetKernelArg(update_buffer_kernel, i + 1, + sizeof(cl_mem), (void *)&(buffers[i])); + } + + if (err != CL_SUCCESS) + { + print_error(err, + "Error: Failed to set arg values for kernel\n"); + goto CLEANUP; + } + err = clEnqueueNDRangeKernel(cmd_queue1, update_buffer_kernel, + 1, NULL, global_work_size, NULL, 0, + NULL, NULL); + if (err != CL_SUCCESS) + { + print_error(err, + "Error: Failed to launch update_buffer_kernel," + " error\n"); + goto CLEANUP; + } + if (iter != (maxIter - 1)) + { + vkWaitForFences(vkDevice, 1, &fence, VK_TRUE, UINT64_MAX); + clFinish(cmd_queue1); + } + } + error_2 = (uint8_t *)malloc(sizeof(uint8_t)); + if (NULL == error_2) + { + log_error("Not able to allocate memory\n"); + goto CLEANUP; + } + + error_1 = clCreateBuffer(context, CL_MEM_WRITE_ONLY, + sizeof(uint8_t), NULL, &err); + if (CL_SUCCESS != err) + { + print_error(err, "Error: clCreateBuffer \n"); + goto CLEANUP; + } + uint8_t val = 0; + err = clEnqueueWriteBuffer(cmd_queue1, error_1, CL_TRUE, 0, + sizeof(uint8_t), &val, 0, NULL, NULL); + if (CL_SUCCESS != err) + { + print_error(err, "Error: clEnqueueWriteBuffer \n"); + goto CLEANUP; + } + + int calc_max_iter = (maxIter * 2); + for (int i = 0; i < vkBufferList.size(); i++) + { + err = clSetKernelArg(verify_kernel, 0, sizeof(cl_mem), + (void *)&(buffers[i])); + 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); + if (err != CL_SUCCESS) + { + print_error( + err, + "Error: Failed to set arg values for verify_kernel \n"); + goto CLEANUP; + } + err = clEnqueueNDRangeKernel(cmd_queue1, verify_kernel, 1, NULL, + global_work_size, NULL, 0, NULL, + NULL); + if (err != CL_SUCCESS) + { + print_error( + err, "Error: Failed to launch verify_kernel, error\n"); + goto CLEANUP; + } + + err = clEnqueueReadBuffer(cmd_queue1, error_1, CL_TRUE, 0, + sizeof(uint8_t), error_2, 0, NULL, + NULL); + if (err != CL_SUCCESS) + { + print_error(err, "Error: Failed read output, error \n"); + goto CLEANUP; + } + if (*error_2 == 1) + { + log_error("&&&& vulkan_opencl_buffer test FAILED\n"); + goto CLEANUP; + } + } + for (size_t i = 0; i < vkBufferList.size(); i++) + { + delete vkBufferListDeviceMemory[i]; + delete externalMemory[i]; + } + vkBufferListDeviceMemory.erase(vkBufferListDeviceMemory.begin(), + vkBufferListDeviceMemory.begin() + + numBuffers); + externalMemory.erase(externalMemory.begin(), + externalMemory.begin() + numBuffers); + } + } +CLEANUP: + for (size_t i = 0; i < vkBufferListDeviceMemory.size(); i++) + { + if (vkBufferListDeviceMemory[i]) + { + delete vkBufferListDeviceMemory[i]; + } + if (externalMemory[i]) + { + delete externalMemory[i]; + } + } + if (error_2) free(error_2); + if (error_1) clReleaseMemObject(error_1); + + vkDestroyFence(vkDevice, fence, nullptr); + + return err; +} + +int run_test_with_multi_import_same_ctx_fence( + 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) +{ + size_t global_work_size[1]; + uint8_t *error_2; + cl_mem error_1; + int numImports = numBuffers; + cl_kernel update_buffer_kernel[MAX_IMPORTS]; + int err = CL_SUCCESS; + int calc_max_iter; + bool withOffset; + uint32_t pBufferSize; + + const std::vector + vkExternalMemoryHandleTypeList = + getSupportedVulkanExternalMemoryHandleTypeList(); + + VulkanQueue &vkQueue = vkDevice.getQueue(); + + std::vector vkBufferShader = readFile("buffer.spv"); + + VulkanShaderModule vkBufferShaderModule(vkDevice, vkBufferShader); + VulkanDescriptorSetLayoutBindingList vkDescriptorSetLayoutBindingList( + MAX_BUFFERS + 1, VULKAN_DESCRIPTOR_TYPE_STORAGE_BUFFER); + VulkanDescriptorSetLayout vkDescriptorSetLayout( + vkDevice, vkDescriptorSetLayoutBindingList); + VulkanPipelineLayout vkPipelineLayout(vkDevice, vkDescriptorSetLayout); + VulkanComputePipeline vkComputePipeline(vkDevice, vkPipelineLayout, + vkBufferShaderModule); + + VulkanDescriptorPool vkDescriptorPool(vkDevice, + vkDescriptorSetLayoutBindingList); + VulkanDescriptorSet vkDescriptorSet(vkDevice, vkDescriptorPool, + vkDescriptorSetLayout); + + const uint32_t maxIter = innerIterations; + VulkanCommandPool vkCommandPool(vkDevice); + VulkanCommandBuffer vkCommandBuffer(vkDevice, vkCommandPool); + + VulkanBuffer vkParamsBuffer(vkDevice, sizeof(Params)); + VulkanDeviceMemory vkParamsDeviceMemory( + vkDevice, vkParamsBuffer.getSize(), + getVulkanMemoryType(vkDevice, + VULKAN_MEMORY_TYPE_PROPERTY_HOST_VISIBLE_COHERENT)); + vkParamsDeviceMemory.bindBuffer(vkParamsBuffer); + std::vector vkBufferListDeviceMemory; + std::vector> externalMemory; + + VkFence fence; + VkFenceCreateInfo fenceInfo{}; + fenceInfo.sType = VK_STRUCTURE_TYPE_FENCE_CREATE_INFO; + fenceInfo.pNext = nullptr; + fenceInfo.flags = 0; + + VkResult vkStatus = vkCreateFence(vkDevice, &fenceInfo, nullptr, &fence); + + if (vkStatus != VK_SUCCESS) + { + print_error(vkStatus, "Error: Failed create fence.\n"); + goto CLEANUP; + } + + for (size_t emhtIdx = 0; emhtIdx < vkExternalMemoryHandleTypeList.size(); + emhtIdx++) + { + VulkanExternalMemoryHandleType vkExternalMemoryHandleType = + vkExternalMemoryHandleTypeList[emhtIdx]; + log_info("External memory handle type: %d\n", + vkExternalMemoryHandleType); + + VulkanBuffer vkDummyBuffer(vkDevice, 4 * 1024, + vkExternalMemoryHandleType); + const VulkanMemoryTypeList &memoryTypeList = + vkDummyBuffer.getMemoryTypeList(); + + for (size_t mtIdx = 0; mtIdx < memoryTypeList.size(); mtIdx++) + { + const VulkanMemoryType &memoryType = memoryTypeList[mtIdx]; + + log_info("Memory type index: %d\n", (uint32_t)memoryType); + log_info("Memory type property: %d\n", + memoryType.getMemoryTypeProperty()); + for (unsigned int withOffset = 0; + withOffset <= (unsigned int)enableOffset; withOffset++) + { + log_info("Running withOffset case %d\n", (uint32_t)withOffset); + if (withOffset) + { + pBufferSize = bufferSizeForOffset; + } + else + { + pBufferSize = bufferSize; + } + cl_mem buffers[MAX_BUFFERS][MAX_IMPORTS]; + VulkanBufferList vkBufferList(numBuffers, vkDevice, pBufferSize, + vkExternalMemoryHandleType); + uint32_t interBufferOffset = + (uint32_t)(vkBufferList[0].getSize()); + + for (size_t bIdx = 0; bIdx < numBuffers; bIdx++) + { + if (withOffset == 0) + { + vkBufferListDeviceMemory.push_back( + new VulkanDeviceMemory(vkDevice, pBufferSize, + memoryType, + vkExternalMemoryHandleType)); + } + if (withOffset == 1) + { + uint32_t totalSize = + (uint32_t)(vkBufferList.size() * interBufferOffset); + vkBufferListDeviceMemory.push_back( + new VulkanDeviceMemory(vkDevice, totalSize, + memoryType, + vkExternalMemoryHandleType)); + } + std::vector pExternalMemory; + for (size_t cl_bIdx = 0; cl_bIdx < numImports; cl_bIdx++) + { + pExternalMemory.push_back(new clExternalMemory( + vkBufferListDeviceMemory[bIdx], + vkExternalMemoryHandleType, + withOffset * bIdx * interBufferOffset, pBufferSize, + context, deviceId)); + } + externalMemory.push_back(pExternalMemory); + } + + clFinish(cmd_queue1); + Params *params = (Params *)vkParamsDeviceMemory.map(); + params->numBuffers = numBuffers; + params->bufferSize = pBufferSize; + params->interBufferOffset = interBufferOffset * withOffset; + vkParamsDeviceMemory.unmap(); + vkDescriptorSet.update(0, vkParamsBuffer); + for (size_t bIdx = 0; bIdx < vkBufferList.size(); bIdx++) + { + size_t buffer_size = vkBufferList[bIdx].getSize(); + vkBufferListDeviceMemory[bIdx]->bindBuffer( + vkBufferList[bIdx], + bIdx * interBufferOffset * withOffset); + for (size_t cl_bIdx = 0; cl_bIdx < numImports; cl_bIdx++) + { + buffers[bIdx][cl_bIdx] = + externalMemory[bIdx][cl_bIdx] + ->getExternalMemoryBuffer(); + } + vkDescriptorSet.update((uint32_t)bIdx + 1, + vkBufferList[bIdx]); + } + vkCommandBuffer.begin(); + vkCommandBuffer.bindPipeline(vkComputePipeline); + vkCommandBuffer.bindDescriptorSets( + vkComputePipeline, vkPipelineLayout, vkDescriptorSet); + vkCommandBuffer.dispatch(512, 1, 1); + vkCommandBuffer.end(); + for (int i = 0; i < numImports; i++) + { + update_buffer_kernel[i] = (numBuffers == 1) + ? kernel[0] + : ((numBuffers == 2) ? kernel[1] : kernel[2]); + } + // global work size should be less than or equal to + // bufferSizeList[i] + global_work_size[0] = pBufferSize; + + for (uint32_t iter = 0; iter < maxIter; iter++) + { + if (iter == 0) + { + vkQueue.submit(vkCommandBuffer, fence); + } + else + { + vkQueue.submit(vkCommandBuffer, fence); + } + vkWaitForFences(vkDevice, 1, &fence, VK_TRUE, UINT64_MAX); + for (uint8_t launchIter = 0; launchIter < numImports; + launchIter++) + { + err = clSetKernelArg(update_buffer_kernel[launchIter], + 0, sizeof(uint32_t), + (void *)&pBufferSize); + for (int i = 0; i < numBuffers; i++) + { + err |= clSetKernelArg( + update_buffer_kernel[launchIter], i + 1, + sizeof(cl_mem), + (void *)&(buffers[i][launchIter])); + } + + if (err != CL_SUCCESS) + { + print_error(err, + "Error: Failed to set arg values for " + "kernel\n "); + goto CLEANUP; + } + err = clEnqueueNDRangeKernel( + cmd_queue1, update_buffer_kernel[launchIter], 1, + NULL, global_work_size, NULL, 0, NULL, NULL); + if (err != CL_SUCCESS) + { + print_error(err, + "Error: Failed to launch " + "update_buffer_kernel, error\n "); + goto CLEANUP; + } + } + if (iter != (maxIter - 1)) + { + vkWaitForFences(vkDevice, 1, &fence, VK_TRUE, + UINT64_MAX); + } + } + error_2 = (uint8_t *)malloc(sizeof(uint8_t)); + if (NULL == error_2) + { + log_error("Not able to allocate memory\n"); + goto CLEANUP; + } + + error_1 = clCreateBuffer(context, CL_MEM_WRITE_ONLY, + sizeof(uint8_t), NULL, &err); + if (CL_SUCCESS != err) + { + print_error(err, "Error: clCreateBuffer \n"); + goto CLEANUP; + } + uint8_t val = 0; + err = + clEnqueueWriteBuffer(cmd_queue1, error_1, CL_TRUE, 0, + sizeof(uint8_t), &val, 0, NULL, NULL); + if (CL_SUCCESS != err) + { + print_error(err, "Error: clEnqueueWriteBuffer \n"); + goto CLEANUP; + } + calc_max_iter = maxIter * (numBuffers + 1); + + for (int i = 0; i < vkBufferList.size(); i++) + { + err = clSetKernelArg(verify_kernel, 0, sizeof(cl_mem), + (void *)&(buffers[i][0])); + err |= clSetKernelArg(verify_kernel, 1, sizeof(int), + &pBufferSize); + err |= clSetKernelArg(verify_kernel, 2, sizeof(int), + &calc_max_iter); + err |= clSetKernelArg(verify_kernel, 3, sizeof(cl_mem), + (void *)&error_1); + if (err != CL_SUCCESS) + { + print_error(err, + "Error: Failed to set arg values for " + "verify_kernel \n"); + goto CLEANUP; + } + err = clEnqueueNDRangeKernel(cmd_queue1, verify_kernel, 1, + NULL, global_work_size, NULL, + 0, NULL, NULL); + if (err != CL_SUCCESS) + { + print_error( + err, + "Error: Failed to launch verify_kernel, error\n"); + goto CLEANUP; + } + + err = clEnqueueReadBuffer(cmd_queue1, error_1, CL_TRUE, 0, + sizeof(uint8_t), error_2, 0, NULL, + NULL); + if (err != CL_SUCCESS) + { + print_error(err, "Error: Failed read output, error \n"); + goto CLEANUP; + } + if (*error_2 == 1) + { + log_error("&&&& vulkan_opencl_buffer test FAILED\n"); + goto CLEANUP; + } + } + 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: + for (size_t i = 0; i < vkBufferListDeviceMemory.size(); i++) + { + if (vkBufferListDeviceMemory[i]) + { + delete vkBufferListDeviceMemory[i]; + } + } + for (size_t i = 0; i < externalMemory.size(); i++) + { + for (size_t j = 0; j < externalMemory[i].size(); j++) + { + if (externalMemory[i][j]) + { + delete externalMemory[i][j]; + } + } + } + if (error_2) free(error_2); + if (error_1) clReleaseMemObject(error_1); + + vkDestroyFence(vkDevice, fence, nullptr); + + return err; +} + +int run_test_with_multi_import_diff_ctx_fence( + 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) +{ + size_t global_work_size[1]; + uint8_t *error_3; + cl_mem error_1; + cl_mem error_2; + int numImports = numBuffers; + cl_kernel update_buffer_kernel1[MAX_IMPORTS]; + cl_kernel update_buffer_kernel2[MAX_IMPORTS]; + + int err = CL_SUCCESS; + int calc_max_iter; + bool withOffset; + uint32_t pBufferSize; + + const std::vector + vkExternalMemoryHandleTypeList = + getSupportedVulkanExternalMemoryHandleTypeList(); + VulkanExternalSemaphoreHandleType vkExternalSemaphoreHandleType = + getSupportedVulkanExternalSemaphoreHandleTypeList()[0]; + + VulkanQueue &vkQueue = vkDevice.getQueue(); + + std::vector vkBufferShader = readFile("buffer.spv"); + + VulkanShaderModule vkBufferShaderModule(vkDevice, vkBufferShader); + VulkanDescriptorSetLayoutBindingList vkDescriptorSetLayoutBindingList( + MAX_BUFFERS + 1, VULKAN_DESCRIPTOR_TYPE_STORAGE_BUFFER); + VulkanDescriptorSetLayout vkDescriptorSetLayout( + vkDevice, vkDescriptorSetLayoutBindingList); + VulkanPipelineLayout vkPipelineLayout(vkDevice, vkDescriptorSetLayout); + VulkanComputePipeline vkComputePipeline(vkDevice, vkPipelineLayout, + vkBufferShaderModule); + + VulkanDescriptorPool vkDescriptorPool(vkDevice, + vkDescriptorSetLayoutBindingList); + VulkanDescriptorSet vkDescriptorSet(vkDevice, vkDescriptorPool, + vkDescriptorSetLayout); + + const uint32_t maxIter = innerIterations; + VulkanCommandPool vkCommandPool(vkDevice); + VulkanCommandBuffer vkCommandBuffer(vkDevice, vkCommandPool); + + VulkanBuffer vkParamsBuffer(vkDevice, sizeof(Params)); + VulkanDeviceMemory vkParamsDeviceMemory( + vkDevice, vkParamsBuffer.getSize(), + getVulkanMemoryType(vkDevice, + VULKAN_MEMORY_TYPE_PROPERTY_HOST_VISIBLE_COHERENT)); + vkParamsDeviceMemory.bindBuffer(vkParamsBuffer); + std::vector vkBufferListDeviceMemory; + std::vector> externalMemory1; + std::vector> externalMemory2; + + VkFence fence; + VkFenceCreateInfo fenceInfo{}; + fenceInfo.sType = VK_STRUCTURE_TYPE_FENCE_CREATE_INFO; + fenceInfo.pNext = nullptr; + fenceInfo.flags = 0; + + VkResult vkStatus = vkCreateFence(vkDevice, &fenceInfo, nullptr, &fence); + if (vkStatus != VK_SUCCESS) + { + print_error(vkStatus, "Error: Failed create fence.\n"); + goto CLEANUP; + } + + for (size_t emhtIdx = 0; emhtIdx < vkExternalMemoryHandleTypeList.size(); + emhtIdx++) + { + VulkanExternalMemoryHandleType vkExternalMemoryHandleType = + vkExternalMemoryHandleTypeList[emhtIdx]; + log_info("External memory handle type:%d\n", + vkExternalMemoryHandleType); + + VulkanBuffer vkDummyBuffer(vkDevice, 4 * 1024, + vkExternalMemoryHandleType); + const VulkanMemoryTypeList &memoryTypeList = + vkDummyBuffer.getMemoryTypeList(); + + for (size_t mtIdx = 0; mtIdx < memoryTypeList.size(); mtIdx++) + { + const VulkanMemoryType &memoryType = memoryTypeList[mtIdx]; + + log_info("Memory type index: %d\n", (uint32_t)memoryType); + log_info("Memory type property: %d\n", + memoryType.getMemoryTypeProperty()); + + for (unsigned int withOffset = 0; + withOffset <= (unsigned int)enableOffset; withOffset++) + { + log_info("Running withOffset case %d\n", (uint32_t)withOffset); + cl_mem buffers1[MAX_BUFFERS][MAX_IMPORTS]; + cl_mem buffers2[MAX_BUFFERS][MAX_IMPORTS]; + if (withOffset) + { + pBufferSize = bufferSizeForOffset; + } + else + { + pBufferSize = bufferSize; + } + VulkanBufferList vkBufferList(numBuffers, vkDevice, pBufferSize, + vkExternalMemoryHandleType); + uint32_t interBufferOffset = + (uint32_t)(vkBufferList[0].getSize()); + + for (size_t bIdx = 0; bIdx < numBuffers; bIdx++) + { + if (withOffset == 0) + { + vkBufferListDeviceMemory.push_back( + new VulkanDeviceMemory(vkDevice, pBufferSize, + memoryType, + vkExternalMemoryHandleType)); + } + if (withOffset == 1) + { + uint32_t totalSize = + (uint32_t)(vkBufferList.size() * interBufferOffset); + vkBufferListDeviceMemory.push_back( + new VulkanDeviceMemory(vkDevice, totalSize, + memoryType, + vkExternalMemoryHandleType)); + } + std::vector pExternalMemory1; + std::vector pExternalMemory2; + for (size_t cl_bIdx = 0; cl_bIdx < numImports; cl_bIdx++) + { + pExternalMemory1.push_back(new clExternalMemory( + vkBufferListDeviceMemory[bIdx], + vkExternalMemoryHandleType, + withOffset * bIdx * interBufferOffset, pBufferSize, + context, deviceId)); + pExternalMemory2.push_back(new clExternalMemory( + vkBufferListDeviceMemory[bIdx], + vkExternalMemoryHandleType, + withOffset * bIdx * interBufferOffset, pBufferSize, + context2, deviceId)); + } + externalMemory1.push_back(pExternalMemory1); + externalMemory2.push_back(pExternalMemory2); + } + + clFinish(cmd_queue1); + Params *params = (Params *)vkParamsDeviceMemory.map(); + params->numBuffers = numBuffers; + params->bufferSize = pBufferSize; + params->interBufferOffset = interBufferOffset * withOffset; + vkParamsDeviceMemory.unmap(); + vkDescriptorSet.update(0, vkParamsBuffer); + for (size_t bIdx = 0; bIdx < vkBufferList.size(); bIdx++) + { + size_t buffer_size = vkBufferList[bIdx].getSize(); + vkBufferListDeviceMemory[bIdx]->bindBuffer( + vkBufferList[bIdx], + bIdx * interBufferOffset * withOffset); + for (size_t cl_bIdx = 0; cl_bIdx < numImports; cl_bIdx++) + { + buffers1[bIdx][cl_bIdx] = + externalMemory1[bIdx][cl_bIdx] + ->getExternalMemoryBuffer(); + buffers2[bIdx][cl_bIdx] = + externalMemory2[bIdx][cl_bIdx] + ->getExternalMemoryBuffer(); + } + vkDescriptorSet.update((uint32_t)bIdx + 1, + vkBufferList[bIdx]); + } + + vkCommandBuffer.begin(); + vkCommandBuffer.bindPipeline(vkComputePipeline); + vkCommandBuffer.bindDescriptorSets( + vkComputePipeline, vkPipelineLayout, vkDescriptorSet); + vkCommandBuffer.dispatch(512, 1, 1); + vkCommandBuffer.end(); + + for (int i = 0; i < numImports; i++) + { + update_buffer_kernel1[i] = (numBuffers == 1) + ? kernel1[0] + : ((numBuffers == 2) ? kernel1[1] : kernel1[2]); + update_buffer_kernel2[i] = (numBuffers == 1) + ? kernel2[0] + : ((numBuffers == 2) ? kernel2[1] : kernel2[2]); + } + + // global work size should be less than or equal + // to bufferSizeList[i] + global_work_size[0] = pBufferSize; + + for (uint32_t iter = 0; iter < maxIter; iter++) + { + if (iter == 0) + { + vkQueue.submit(vkCommandBuffer, fence); + } + else + { + vkQueue.submit(vkCommandBuffer, fence); + } + vkWaitForFences(vkDevice, 1, &fence, VK_TRUE, UINT64_MAX); + + for (uint8_t launchIter = 0; launchIter < numImports; + launchIter++) + { + err = clSetKernelArg(update_buffer_kernel1[launchIter], + 0, sizeof(uint32_t), + (void *)&pBufferSize); + for (int i = 0; i < numBuffers; i++) + { + err |= clSetKernelArg( + update_buffer_kernel1[launchIter], i + 1, + sizeof(cl_mem), + (void *)&(buffers1[i][launchIter])); + } + + if (err != CL_SUCCESS) + { + print_error(err, + "Error: Failed to set arg values for " + "kernel\n "); + goto CLEANUP; + } + err = clEnqueueNDRangeKernel( + cmd_queue1, update_buffer_kernel1[launchIter], 1, + NULL, global_work_size, NULL, 0, NULL, NULL); + if (err != CL_SUCCESS) + { + print_error(err, + "Error: Failed to launch " + "update_buffer_kernel, error\n"); + goto CLEANUP; + } + } + if (iter != (maxIter - 1)) + { + vkWaitForFences(vkDevice, 1, &fence, VK_TRUE, + UINT64_MAX); + } + } + clFinish(cmd_queue1); + for (uint32_t iter = 0; iter < maxIter; iter++) + { + if (iter == 0) + { + vkQueue.submit(vkCommandBuffer, fence); + } + else + { + vkQueue.submit(vkCommandBuffer, fence); + } + vkWaitForFences(vkDevice, 1, &fence, VK_TRUE, UINT64_MAX); + + for (uint8_t launchIter = 0; launchIter < numImports; + launchIter++) + { + err = clSetKernelArg(update_buffer_kernel2[launchIter], + 0, sizeof(uint32_t), + (void *)&bufferSize); + for (int i = 0; i < numBuffers; i++) + { + err |= clSetKernelArg( + update_buffer_kernel2[launchIter], i + 1, + sizeof(cl_mem), + (void *)&(buffers2[i][launchIter])); + } + + if (err != CL_SUCCESS) + { + print_error(err, + "Error: Failed to set arg values for " + "kernel\n "); + goto CLEANUP; + } + err = clEnqueueNDRangeKernel( + cmd_queue2, update_buffer_kernel2[launchIter], 1, + NULL, global_work_size, NULL, 0, NULL, NULL); + if (err != CL_SUCCESS) + { + print_error(err, + "Error: Failed to launch " + "update_buffer_kernel, error\n "); + goto CLEANUP; + } + } + if (iter != (maxIter - 1)) + { + vkWaitForFences(vkDevice, 1, &fence, VK_TRUE, + UINT64_MAX); + } + } + clFinish(cmd_queue2); + error_3 = (uint8_t *)malloc(sizeof(uint8_t)); + if (NULL == error_3) + { + log_error("Not able to allocate memory\n"); + goto CLEANUP; + } + + error_1 = clCreateBuffer(context, CL_MEM_WRITE_ONLY, + sizeof(uint8_t), NULL, &err); + if (CL_SUCCESS != err) + { + print_error(err, "Error: clCreateBuffer \n"); + goto CLEANUP; + } + error_2 = clCreateBuffer(context2, CL_MEM_WRITE_ONLY, + sizeof(uint8_t), NULL, &err); + if (CL_SUCCESS != err) + { + print_error(err, "Error: clCreateBuffer \n"); + goto CLEANUP; + } + uint8_t val = 0; + err = + clEnqueueWriteBuffer(cmd_queue1, error_1, CL_TRUE, 0, + sizeof(uint8_t), &val, 0, NULL, NULL); + if (err != CL_SUCCESS) + { + print_error(err, "Error: Failed read output, error \n"); + goto CLEANUP; + } + + err = + clEnqueueWriteBuffer(cmd_queue2, error_2, CL_TRUE, 0, + sizeof(uint8_t), &val, 0, NULL, NULL); + if (err != CL_SUCCESS) + { + print_error(err, "Error: Failed read output, error \n"); + goto CLEANUP; + } + + 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); + if (err != CL_SUCCESS) + { + print_error(err, + "Error: Failed to set arg values for " + "verify_kernel \n"); + goto CLEANUP; + } + err = clEnqueueNDRangeKernel(cmd_queue1, verify_kernel, 1, + NULL, global_work_size, NULL, + 0, NULL, NULL); + if (err != CL_SUCCESS) + { + print_error(err, + "Error: Failed to launch verify_kernel," + "error\n"); + goto CLEANUP; + } + + err = clEnqueueReadBuffer(cmd_queue1, error_1, CL_TRUE, 0, + sizeof(uint8_t), error_3, 0, NULL, + NULL); + if (err != CL_SUCCESS) + { + print_error(err, "Error: Failed read output, error\n"); + goto CLEANUP; + } + if (*error_3 == 1) + { + log_error("&&&& vulkan_opencl_buffer test FAILED\n"); + goto CLEANUP; + } + } + *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); + if (err != CL_SUCCESS) + { + print_error(err, + "Error: Failed to set arg values for " + "verify_kernel \n"); + goto CLEANUP; + } + err = clEnqueueNDRangeKernel(cmd_queue2, verify_kernel2, 1, + NULL, global_work_size, NULL, + 0, NULL, NULL); + if (err != CL_SUCCESS) + { + print_error(err, + "Error: Failed to launch verify_kernel," + "error\n"); + goto CLEANUP; + } + + err = clEnqueueReadBuffer(cmd_queue2, error_2, CL_TRUE, 0, + sizeof(uint8_t), error_3, 0, NULL, + NULL); + if (err != CL_SUCCESS) + { + print_error(err, "Error: Failed read output, error\n"); + goto CLEANUP; + } + if (*error_3 == 1) + { + log_error("&&&& vulkan_opencl_buffer test FAILED\n"); + goto CLEANUP; + } + } + 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(); + } + } + } +CLEANUP: + for (size_t i = 0; i < vkBufferListDeviceMemory.size(); i++) + { + if (vkBufferListDeviceMemory[i]) + { + delete vkBufferListDeviceMemory[i]; + } + } + for (size_t i = 0; i < externalMemory1.size(); i++) + { + for (size_t j = 0; j < externalMemory1[i].size(); j++) + { + if (externalMemory1[i][j]) + { + delete externalMemory1[i][j]; + } + } + } + for (size_t i = 0; i < externalMemory2.size(); i++) + { + for (size_t j = 0; j < externalMemory2[i].size(); j++) + { + if (externalMemory2[i][j]) + { + delete externalMemory2[i][j]; + } + } + } + if (error_3) free(error_3); + if (error_1) clReleaseMemObject(error_1); + if (error_2) clReleaseMemObject(error_2); + + vkDestroyFence(vkDevice, fence, nullptr); + + return err; +} + +int test_buffer_common_fence(cl_device_id device_, cl_context context_, + cl_command_queue queue_, int numElements_) +{ + + 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 }; + errNum = clGetPlatformIDs(1, &platform, NULL); + if (errNum != CL_SUCCESS) + { + print_error(errNum, "Error: Failed to get platform\n"); + goto CLEANUP; + } + + errNum = + clGetDeviceIDs(platform, CL_DEVICE_TYPE_GPU, 0, NULL, &num_devices); + if (CL_SUCCESS != errNum) + { + print_error(errNum, "clGetDeviceIDs failed in returning of devices\n"); + goto CLEANUP; + } + devices = (cl_device_id *)malloc(num_devices * sizeof(cl_device_id)); + if (NULL == devices) + { + errNum = CL_OUT_OF_HOST_MEMORY; + print_error(errNum, "Unable to allocate memory for devices\n"); + goto CLEANUP; + } + errNum = clGetDeviceIDs(platform, CL_DEVICE_TYPE_GPU, num_devices, devices, + NULL); + if (CL_SUCCESS != errNum) + { + print_error(errNum, "Failed to get deviceID.\n"); + goto CLEANUP; + } + contextProperties[1] = (cl_context_properties)platform; + log_info("Assigned contextproperties for platform\n"); + for (device_no = 0; device_no < num_devices; device_no++) + { + errNum = clGetDeviceInfo(devices[device_no], CL_DEVICE_EXTENSIONS, 0, + NULL, &extensionSize); + if (CL_SUCCESS != errNum) + { + print_error(errNum, + "Error in clGetDeviceInfo for getting device_extension " + "size....\n"); + goto CLEANUP; + } + extensions = (char *)malloc(extensionSize); + if (NULL == extensions) + { + print_error(errNum, "Unable to allocate memory for extensions\n"); + errNum = CL_OUT_OF_HOST_MEMORY; + goto CLEANUP; + } + errNum = clGetDeviceInfo(devices[device_no], CL_DEVICE_EXTENSIONS, + extensionSize, extensions, NULL); + if (CL_SUCCESS != errNum) + { + print_error(errNum, + "Error in clGetDeviceInfo for device_extension\n"); + goto CLEANUP; + } + errNum = clGetDeviceInfo(devices[device_no], CL_DEVICE_UUID_KHR, + CL_UUID_SIZE_KHR, uuid, &extensionSize); + if (CL_SUCCESS != errNum) + { + print_error(errNum, "clGetDeviceInfo failed\n"); + goto CLEANUP; + } + errNum = + memcmp(uuid, vkDevice.getPhysicalDevice().getUUID(), VK_UUID_SIZE); + if (errNum == 0) + { + break; + } + } + if (device_no >= num_devices) + { + errNum = EXIT_FAILURE; + print_error(errNum, + "OpenCL error: " + "No Vulkan-OpenCL Interop capable GPU found.\n"); + goto CLEANUP; + } + deviceId = devices[device_no]; + context = clCreateContextFromType(contextProperties, CL_DEVICE_TYPE_GPU, + NULL, NULL, &errNum); + if (CL_SUCCESS != errNum) + { + print_error(errNum, "error creating context\n"); + goto CLEANUP; + } + log_info("Successfully created context !!!\n"); + + cmd_queue1 = clCreateCommandQueue(context, devices[device_no], 0, &errNum); + if (CL_SUCCESS != errNum) + { + errNum = CL_INVALID_COMMAND_QUEUE; + print_error(errNum, "Error: Failed to create command queue!\n"); + goto CLEANUP; + } + cmd_queue2 = clCreateCommandQueue(context, devices[device_no], 0, &errNum); + if (CL_SUCCESS != errNum) + { + errNum = CL_INVALID_COMMAND_QUEUE; + print_error(errNum, "Error: Failed to create command queue!\n"); + goto CLEANUP; + } + 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); + if (errNum != CL_SUCCESS) + { + print_error(errNum, "Error: Failed to build program \n"); + return errNum; + } + // create the kernel + kernel[i] = clCreateKernel(program[i], "clUpdateBuffer", &errNum); + if (errNum != CL_SUCCESS) + { + print_error(errNum, "clCreateKernel failed \n"); + return errNum; + } + } + + 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); + if (errNum != CL_SUCCESS) + { + log_error("Error: Failed to build program2\n"); + return errNum; + } + verify_kernel = clCreateKernel(program_verify, "checkKernel", &errNum); + if (errNum != CL_SUCCESS) + { + print_error(errNum, "clCreateKernel failed \n"); + return errNum; + } + + if (multiCtx) // different context guard + { + context2 = clCreateContextFromType( + contextProperties, CL_DEVICE_TYPE_GPU, NULL, NULL, &errNum); + if (CL_SUCCESS != errNum) + { + print_error(errNum, "error creating context\n"); + goto CLEANUP; + } + cmd_queue3 = + clCreateCommandQueue(context2, devices[device_no], 0, &errNum); + if (CL_SUCCESS != errNum) + { + errNum = CL_INVALID_COMMAND_QUEUE; + print_error(errNum, "Error: Failed to create command queue!\n"); + goto CLEANUP; + } + for (int i = 0; i < 3; i++) + { + 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); + if (errNum != CL_SUCCESS) + { + print_error(errNum, "Error: Failed to build program \n"); + return errNum; + } + // create the kernel + kernel2[i] = clCreateKernel(program[i], "clUpdateBuffer", &errNum); + if (errNum != CL_SUCCESS) + { + print_error(errNum, "clCreateKernel failed \n"); + return errNum; + } + } + program_source_length = strlen(program_source_const_verify); + program_verify = + clCreateProgramWithSource(context2, 1, &program_source_const_verify, + &program_source_length, &errNum); + errNum = clBuildProgram(program_verify, 0, NULL, NULL, NULL, NULL); + if (errNum != CL_SUCCESS) + { + log_error("Error: Failed to build program2\n"); + return errNum; + } + verify_kernel2 = clCreateKernel(program_verify, "checkKernel", &errNum); + if (errNum != CL_SUCCESS) + { + print_error(errNum, "clCreateKernel failed \n"); + return errNum; + } + } + + 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++) + { + uint32_t bufferSize = bufferSizeList[sizeIdx]; + uint32_t bufferSizeForOffset = bufferSizeListforOffset[sizeIdx]; + log_info("&&&& RUNNING vulkan_opencl_buffer test for Buffer size: " + "%d\n", + bufferSize); + if (multiImport && !multiCtx) + { + errNum = run_test_with_multi_import_same_ctx_fence( + context, cmd_queue1, kernel, verify_kernel, vkDevice, + numBuffers, bufferSize, bufferSizeForOffset); + } + else if (multiImport && multiCtx) + { + errNum = run_test_with_multi_import_diff_ctx_fence( + context, context2, cmd_queue1, cmd_queue3, kernel, kernel2, + verify_kernel, verify_kernel2, vkDevice, numBuffers, + bufferSize, bufferSizeForOffset); + } + else if (numCQ == 2) + { + errNum = run_test_with_two_queue_fence( + context, cmd_queue1, cmd_queue2, kernel, verify_kernel, + vkDevice, numBuffers + 1, bufferSize); + } + else + { + errNum = run_test_with_one_queue_fence( + context, cmd_queue1, kernel, verify_kernel, vkDevice, + numBuffers, bufferSize); + } + if (errNum != CL_SUCCESS) + { + print_error(errNum, "func_name failed \n"); + goto CLEANUP; + } + } + } + +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); + + if (devices) free(devices); + if (extensions) free(extensions); + + return errNum; +} diff --git a/test_conformance/vulkan/vulkan_interop_common/vulkan_wrapper.cpp b/test_conformance/vulkan/vulkan_interop_common/vulkan_wrapper.cpp index 3ce4af6b0..8d10b9c42 100644 --- a/test_conformance/vulkan/vulkan_interop_common/vulkan_wrapper.cpp +++ b/test_conformance/vulkan/vulkan_interop_common/vulkan_wrapper.cpp @@ -615,6 +615,16 @@ VulkanQueue::VulkanQueue(VkQueue vkQueue): m_vkQueue(vkQueue) {} VulkanQueue::~VulkanQueue() {} +void VulkanQueue::submit(const VulkanCommandBuffer &commandBuffer, const VkFence &fence) +{ + VkSubmitInfo vkSubmitInfo = {}; + vkSubmitInfo.sType = VK_STRUCTURE_TYPE_SUBMIT_INFO; + vkSubmitInfo.pNext = NULL; + vkSubmitInfo.waitSemaphoreCount = (uint32_t)0; + + vkQueueSubmit(m_vkQueue, 1, &vkSubmitInfo, fence); +} + void VulkanQueue::submit(const VulkanSemaphoreList &waitSemaphoreList, const VulkanCommandBufferList &commandBufferList, const VulkanSemaphoreList &signalSemaphoreList) diff --git a/test_conformance/vulkan/vulkan_interop_common/vulkan_wrapper.hpp b/test_conformance/vulkan/vulkan_interop_common/vulkan_wrapper.hpp index 37925ee4a..b22023e5f 100644 --- a/test_conformance/vulkan/vulkan_interop_common/vulkan_wrapper.hpp +++ b/test_conformance/vulkan/vulkan_interop_common/vulkan_wrapper.hpp @@ -157,6 +157,7 @@ class VulkanQueue { public: const VulkanQueueFamily &getQueueFamily(); + void submit(const VulkanCommandBuffer &commandBuffer, const VkFence &fence); void submit(const VulkanSemaphoreList &waitSemaphoreList, const VulkanCommandBufferList &commandBufferList, const VulkanSemaphoreList &signalSemaphoreList); From 443167e4434ce70987c40d0d90c1f32a8acd2265 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Pawe=C5=82=20Jastrz=C4=99bski?= Date: Tue, 21 Feb 2023 12:09:35 +0100 Subject: [PATCH 02/14] Fix clang-format MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Signed-off-by: Paweł Jastrzębski --- .../vulkan/vulkan_interop_common/vulkan_wrapper.cpp | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/test_conformance/vulkan/vulkan_interop_common/vulkan_wrapper.cpp b/test_conformance/vulkan/vulkan_interop_common/vulkan_wrapper.cpp index 8d10b9c42..53a68fd48 100644 --- a/test_conformance/vulkan/vulkan_interop_common/vulkan_wrapper.cpp +++ b/test_conformance/vulkan/vulkan_interop_common/vulkan_wrapper.cpp @@ -615,7 +615,8 @@ VulkanQueue::VulkanQueue(VkQueue vkQueue): m_vkQueue(vkQueue) {} VulkanQueue::~VulkanQueue() {} -void VulkanQueue::submit(const VulkanCommandBuffer &commandBuffer, const VkFence &fence) +void VulkanQueue::submit(const VulkanCommandBuffer &commandBuffer, + const VkFence &fence) { VkSubmitInfo vkSubmitInfo = {}; vkSubmitInfo.sType = VK_STRUCTURE_TYPE_SUBMIT_INFO; From 5faf2dcc3291ee3bdbb8919fac79f35d7a354fe0 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Pawe=C5=82=20Jastrz=C4=99bski?= Date: Fri, 17 Mar 2023 14:05:39 +0100 Subject: [PATCH 03/14] Apply changes for review. MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Apply changes for review: - Make VkFence + clFinish a synchronization option to existing tests instead of creating a separate test that uses fence. Signed-off-by: Paweł Jastrzębski --- test_conformance/vulkan/CMakeLists.txt | 1 - test_conformance/vulkan/main.cpp | 25 +- .../vulkan/test_vulkan_interop_buffer.cpp | 418 +++- .../test_vulkan_interop_buffer_fence.cpp | 1798 ----------------- 4 files changed, 361 insertions(+), 1881 deletions(-) delete mode 100644 test_conformance/vulkan/test_vulkan_interop_buffer_fence.cpp diff --git a/test_conformance/vulkan/CMakeLists.txt b/test_conformance/vulkan/CMakeLists.txt index 004a2c0a8..4f43172af 100644 --- a/test_conformance/vulkan/CMakeLists.txt +++ b/test_conformance/vulkan/CMakeLists.txt @@ -23,7 +23,6 @@ include_directories (${CLConform_INCLUDE_DIR}) set (${MODULE_NAME}_SOURCES main.cpp test_vulkan_interop_buffer.cpp - test_vulkan_interop_buffer_fence.cpp test_vulkan_interop_image.cpp test_vulkan_api_consistency.cpp test_vulkan_platform_device_info.cpp diff --git a/test_conformance/vulkan/main.cpp b/test_conformance/vulkan/main.cpp index f8dad951a..46a96ecd1 100644 --- a/test_conformance/vulkan/main.cpp +++ b/test_conformance/vulkan/main.cpp @@ -52,18 +52,18 @@ static void params_reset() } extern int test_buffer_common(cl_device_id device_, cl_context context_, - cl_command_queue queue_, int numElements_); -extern int test_buffer_common_fence(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_); int test_buffer_single_queue(cl_device_id device_, cl_context context_, - cl_command_queue queue_, int numElements_) + cl_command_queue queue_, int numElements_, + bool fence) { 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_) @@ -71,7 +71,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_) @@ -80,7 +80,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_) @@ -90,14 +90,14 @@ 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_fence(device_, context_, queue_, numElements_); + 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_) @@ -105,7 +105,7 @@ int test_buffer_multiple_queue_fence(cl_device_id device_, cl_context context_, params_reset(); numCQ = 2; log_info("RUNNING TEST WITH TWO QUEUE...... \n\n"); - return test_buffer_common_fence(device_, context_, queue_, numElements_); + return test_buffer_common(device_, context_, queue_, numElements_, true); } int test_buffer_multiImport_sameCtx_fence(cl_device_id device_, cl_context context_, @@ -116,7 +116,7 @@ int test_buffer_multiImport_sameCtx_fence(cl_device_id device_, multiImport = true; log_info("RUNNING TEST WITH MULTIPLE DEVICE MEMORY IMPORT " "IN SAME CONTEXT...... \n\n"); - return test_buffer_common_fence(device_, context_, queue_, numElements_); + return test_buffer_common(device_, context_, queue_, numElements_, true); } int test_buffer_multiImport_diffCtx_fence(cl_device_id device_, cl_context context_, @@ -128,7 +128,7 @@ int test_buffer_multiImport_diffCtx_fence(cl_device_id device_, multiCtx = true; log_info("RUNNING TEST WITH MULTIPLE DEVICE MEMORY IMPORT " "IN DIFFERENT CONTEXT...... \n\n"); - return test_buffer_common_fence(device_, context_, queue_, numElements_); + 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_) @@ -148,6 +148,7 @@ int test_image_multiple_queue(cl_device_id device_, cl_context context_, test_definition test_list[] = { + ADD_TEST(buffer_single_queue), ADD_TEST(buffer_multiple_queue), ADD_TEST(buffer_multiImport_sameCtx), diff --git a/test_conformance/vulkan/test_vulkan_interop_buffer.cpp b/test_conformance/vulkan/test_vulkan_interop_buffer.cpp index 7d3062b3f..488a904b7 100644 --- a/test_conformance/vulkan/test_vulkan_interop_buffer.cpp +++ b/test_conformance/vulkan/test_vulkan_interop_buffer.cpp @@ -39,7 +39,7 @@ struct Params }; } -static const char *kernel_text_numbuffer_1 = " \ +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\ @@ -47,7 +47,7 @@ __kernel void clUpdateBuffer(int bufferSize, __global unsigned char *a) { \n\ } \n\ }"; -static const char *kernel_text_numbuffer_2 = " \ +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\ @@ -56,7 +56,7 @@ __kernel void clUpdateBuffer(int bufferSize, __global unsigned char *a, __global } \n\ }"; -static const char *kernel_text_numbuffer_4 = " \ +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\ @@ -68,7 +68,7 @@ __kernel void clUpdateBuffer(int bufferSize, __global unsigned char *a, __global }"; -static const char *kernel_text_verify = " \ +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\ @@ -82,7 +82,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 +118,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); + VkFence fence; VulkanQueue &vkQueue = vkDevice.getQueue(); @@ -136,10 +138,30 @@ 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) + { + VkFenceCreateInfo fenceInfo{}; + fenceInfo.sType = VK_STRUCTURE_TYPE_FENCE_CREATE_INFO; + fenceInfo.pNext = nullptr; + fenceInfo.flags = 0; + + VkResult vkStatus = + vkCreateFence(vkDevice, &fenceInfo, nullptr, &fence); + + if (vkStatus != VK_SUCCESS) + { + print_error(vkStatus, "Error: Failed create fence.\n"); + vkDestroyFence(vkDevice, fence, nullptr); + return TEST_FAIL; + } + } + else + { + clVk2CLExternalSemaphore = new clExternalSemaphore( + vkVk2CLSemaphore, context, vkExternalSemaphoreHandleType, deviceId); + clCl2VkExternalSemaphore = new clExternalSemaphore( + vkCl2VkSemaphore, context, vkExternalSemaphoreHandleType, deviceId); + } const uint32_t maxIter = innerIterations; VulkanCommandPool vkCommandPool(vkDevice); @@ -229,14 +251,37 @@ int run_test_with_two_queue(cl_context &context, cl_command_queue &cmd_queue1, if (iter == 0) { - vkQueue.submit(vkCommandBuffer, vkVk2CLSemaphore); + if (use_fence) + { + vkQueue.submit(vkCommandBuffer, fence); + } + else + { + vkQueue.submit(vkCommandBuffer, vkVk2CLSemaphore); + } + } + else + { + if (use_fence) + { + vkQueue.submit(vkCommandBuffer, fence); + } + else + { + vkQueue.submit(vkCl2VkSemaphore, vkCommandBuffer, + vkVk2CLSemaphore); + } + } + + if (use_fence) + { + vkWaitForFences(vkDevice, 1, &fence, VK_TRUE, UINT64_MAX); } 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); @@ -288,7 +333,15 @@ int run_test_with_two_queue(cl_context &context, cl_command_queue &cmd_queue1, if (iter != (maxIter - 1)) { - clCl2VkExternalSemaphore->signal(cmd_queue2); + if (use_fence) + { + vkWaitForFences(vkDevice, 1, &fence, VK_TRUE, + UINT64_MAX); + } + else + { + clCl2VkExternalSemaphore->signal(cmd_queue2); + } } } error_2 = (uint8_t *)malloc(sizeof(uint8_t)); @@ -387,8 +440,15 @@ 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) + { + vkDestroyFence(vkDevice, fence, nullptr); + } + else + { + if (clVk2CLExternalSemaphore) delete clVk2CLExternalSemaphore; + if (clCl2VkExternalSemaphore) delete clCl2VkExternalSemaphore; + } if (error_2) free(error_2); if (error_1) clReleaseMemObject(error_1); @@ -398,7 +458,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 +476,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); + VkFence fence; VulkanQueue &vkQueue = vkDevice.getQueue(); @@ -434,10 +495,31 @@ 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) + { + VkFenceCreateInfo fenceInfo{}; + fenceInfo.sType = VK_STRUCTURE_TYPE_FENCE_CREATE_INFO; + fenceInfo.pNext = nullptr; + fenceInfo.flags = 0; + + VkResult vkStatus = + vkCreateFence(vkDevice, &fenceInfo, nullptr, &fence); + + if (vkStatus != VK_SUCCESS) + { + print_error(vkStatus, "Error: Failed create fence.\n"); + vkDestroyFence(vkDevice, fence, nullptr); + return TEST_FAIL; + } + } + 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); @@ -528,14 +610,36 @@ int run_test_with_one_queue(cl_context &context, cl_command_queue &cmd_queue1, { if (iter == 0) { - vkQueue.submit(vkCommandBuffer, vkVk2CLSemaphore); + if (use_fence) + { + vkQueue.submit(vkCommandBuffer, fence); + } + else + { + vkQueue.submit(vkCommandBuffer, vkVk2CLSemaphore); + } } else { - vkQueue.submit(vkCl2VkSemaphore, vkCommandBuffer, - vkVk2CLSemaphore); + if (use_fence) + { + vkQueue.submit(vkCommandBuffer, fence); + } + else + { + vkQueue.submit(vkCl2VkSemaphore, vkCommandBuffer, + vkVk2CLSemaphore); + } + } + + if (use_fence) + { + vkWaitForFences(vkDevice, 1, &fence, VK_TRUE, UINT64_MAX); + } + else + { + clVk2CLExternalSemaphore->wait(cmd_queue1); } - clVk2CLExternalSemaphore->wait(cmd_queue1); err = clSetKernelArg(update_buffer_kernel, 0, sizeof(uint32_t), (void *)&bufferSize); @@ -564,7 +668,16 @@ int run_test_with_one_queue(cl_context &context, cl_command_queue &cmd_queue1, } if (iter != (maxIter - 1)) { - clCl2VkExternalSemaphore->signal(cmd_queue1); + if (use_fence) + { + vkWaitForFences(vkDevice, 1, &fence, VK_TRUE, + UINT64_MAX); + clFinish(cmd_queue1); + } + else + { + clCl2VkExternalSemaphore->signal(cmd_queue1); + } } } error_2 = (uint8_t *)malloc(sizeof(uint8_t)); @@ -656,8 +769,17 @@ 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) + { + vkDestroyFence(vkDevice, fence, nullptr); + } + else + { + if (clVk2CLExternalSemaphore) delete clVk2CLExternalSemaphore; + if (clCl2VkExternalSemaphore) delete clCl2VkExternalSemaphore; + } + if (error_2) free(error_2); if (error_1) clReleaseMemObject(error_1); return err; @@ -666,7 +788,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 +809,7 @@ int run_test_with_multi_import_same_ctx( getSupportedVulkanExternalSemaphoreHandleTypeList()[0]; VulkanSemaphore vkVk2CLSemaphore(vkDevice, vkExternalSemaphoreHandleType); VulkanSemaphore vkCl2VkSemaphore(vkDevice, vkExternalSemaphoreHandleType); + VkFence fence; VulkanQueue &vkQueue = vkDevice.getQueue(); @@ -706,10 +829,31 @@ 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) + { + VkFenceCreateInfo fenceInfo{}; + fenceInfo.sType = VK_STRUCTURE_TYPE_FENCE_CREATE_INFO; + fenceInfo.pNext = nullptr; + fenceInfo.flags = 0; + + VkResult vkStatus = + vkCreateFence(vkDevice, &fenceInfo, nullptr, &fence); + + if (vkStatus != VK_SUCCESS) + { + print_error(vkStatus, "Error: Failed create fence.\n"); + vkDestroyFence(vkDevice, fence, nullptr); + return TEST_FAIL; + } + } + 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); @@ -834,14 +978,38 @@ int run_test_with_multi_import_same_ctx( { if (iter == 0) { - vkQueue.submit(vkCommandBuffer, vkVk2CLSemaphore); + if (use_fence) + { + vkQueue.submit(vkCommandBuffer, fence); + } + else + { + vkQueue.submit(vkCommandBuffer, vkVk2CLSemaphore); + } } else { - vkQueue.submit(vkCl2VkSemaphore, vkCommandBuffer, - vkVk2CLSemaphore); + if (use_fence) + { + vkQueue.submit(vkCommandBuffer, fence); + } + else + { + vkQueue.submit(vkCl2VkSemaphore, vkCommandBuffer, + vkVk2CLSemaphore); + } } - clVk2CLExternalSemaphore->wait(cmd_queue1); + + if (use_fence) + { + vkWaitForFences(vkDevice, 1, &fence, VK_TRUE, + UINT64_MAX); + } + else + { + clVk2CLExternalSemaphore->wait(cmd_queue1); + } + for (uint8_t launchIter = 0; launchIter < numImports; launchIter++) { @@ -876,7 +1044,15 @@ int run_test_with_multi_import_same_ctx( } if (iter != (maxIter - 1)) { - clCl2VkExternalSemaphore->signal(cmd_queue1); + if (use_fence) + { + vkWaitForFences(vkDevice, 1, &fence, VK_TRUE, + UINT64_MAX); + } + else + { + clCl2VkExternalSemaphore->signal(cmd_queue1); + } } } error_2 = (uint8_t *)malloc(sizeof(uint8_t)); @@ -987,8 +1163,17 @@ int run_test_with_multi_import_same_ctx( } } } - if (clVk2CLExternalSemaphore) delete clVk2CLExternalSemaphore; - if (clCl2VkExternalSemaphore) delete clCl2VkExternalSemaphore; + + if (use_fence) + { + vkDestroyFence(vkDevice, fence, nullptr); + } + else + { + if (clVk2CLExternalSemaphore) delete clVk2CLExternalSemaphore; + if (clCl2VkExternalSemaphore) delete clCl2VkExternalSemaphore; + } + if (error_2) free(error_2); if (error_1) clReleaseMemObject(error_1); return err; @@ -998,7 +1183,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 +1209,7 @@ int run_test_with_multi_import_diff_ctx( getSupportedVulkanExternalSemaphoreHandleTypeList()[0]; VulkanSemaphore vkVk2CLSemaphore(vkDevice, vkExternalSemaphoreHandleType); VulkanSemaphore vkCl2VkSemaphore(vkDevice, vkExternalSemaphoreHandleType); + VkFence fence; VulkanQueue &vkQueue = vkDevice.getQueue(); @@ -1042,15 +1229,37 @@ 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); + if (use_fence) + { + VkFenceCreateInfo fenceInfo{}; + fenceInfo.sType = VK_STRUCTURE_TYPE_FENCE_CREATE_INFO; + fenceInfo.pNext = nullptr; + fenceInfo.flags = 0; + + VkResult vkStatus = + vkCreateFence(vkDevice, &fenceInfo, nullptr, &fence); - clVk2CLExternalSemaphore2 = new clExternalSemaphore( - vkVk2CLSemaphore, context2, vkExternalSemaphoreHandleType, deviceId); - clCl2VkExternalSemaphore2 = new clExternalSemaphore( - vkCl2VkSemaphore, context2, vkExternalSemaphoreHandleType, deviceId); + if (vkStatus != VK_SUCCESS) + { + print_error(vkStatus, "Error: Failed create fence.\n"); + vkDestroyFence(vkDevice, fence, nullptr); + return TEST_FAIL; + } + } + 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); @@ -1194,14 +1403,37 @@ int run_test_with_multi_import_diff_ctx( { if (iter == 0) { - vkQueue.submit(vkCommandBuffer, vkVk2CLSemaphore); + if (use_fence) + { + vkQueue.submit(vkCommandBuffer, fence); + } + else + { + vkQueue.submit(vkCommandBuffer, vkVk2CLSemaphore); + } } else { - vkQueue.submit(vkCl2VkSemaphore, vkCommandBuffer, - vkVk2CLSemaphore); + if (use_fence) + { + vkQueue.submit(vkCommandBuffer, fence); + } + else + { + vkQueue.submit(vkCl2VkSemaphore, vkCommandBuffer, + vkVk2CLSemaphore); + } + } + + if (use_fence) + { + vkWaitForFences(vkDevice, 1, &fence, VK_TRUE, + UINT64_MAX); + } + else + { + clVk2CLExternalSemaphore->wait(cmd_queue1); } - clVk2CLExternalSemaphore->wait(cmd_queue1); for (uint8_t launchIter = 0; launchIter < numImports; launchIter++) @@ -1237,7 +1469,15 @@ int run_test_with_multi_import_diff_ctx( } if (iter != (maxIter - 1)) { - clCl2VkExternalSemaphore->signal(cmd_queue1); + if (use_fence) + { + vkWaitForFences(vkDevice, 1, &fence, VK_TRUE, + UINT64_MAX); + } + else + { + clCl2VkExternalSemaphore->signal(cmd_queue1); + } } } clFinish(cmd_queue1); @@ -1245,14 +1485,36 @@ int run_test_with_multi_import_diff_ctx( { if (iter == 0) { - vkQueue.submit(vkCommandBuffer, vkVk2CLSemaphore); + if (use_fence) + { + vkQueue.submit(vkCommandBuffer, fence); + } + else + { + vkQueue.submit(vkCommandBuffer, vkVk2CLSemaphore); + } } else { - vkQueue.submit(vkCl2VkSemaphore, vkCommandBuffer, - vkVk2CLSemaphore); + if (use_fence) + { + vkQueue.submit(vkCommandBuffer, fence); + } + else + { + vkQueue.submit(vkCl2VkSemaphore, vkCommandBuffer, + vkVk2CLSemaphore); + } + } + + if (use_fence) + { + vkQueue.submit(vkCommandBuffer, fence); + } + else + { + clVk2CLExternalSemaphore2->wait(cmd_queue2); } - clVk2CLExternalSemaphore2->wait(cmd_queue2); for (uint8_t launchIter = 0; launchIter < numImports; launchIter++) @@ -1288,7 +1550,15 @@ int run_test_with_multi_import_diff_ctx( } if (iter != (maxIter - 1)) { - clCl2VkExternalSemaphore2->signal(cmd_queue2); + if (use_fence) + { + vkWaitForFences(vkDevice, 1, &fence, VK_TRUE, + UINT64_MAX); + } + else + { + clCl2VkExternalSemaphore2->signal(cmd_queue2); + } } } clFinish(cmd_queue2); @@ -1474,10 +1744,17 @@ 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) + {} + else + { + 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 +1762,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 +2016,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) { diff --git a/test_conformance/vulkan/test_vulkan_interop_buffer_fence.cpp b/test_conformance/vulkan/test_vulkan_interop_buffer_fence.cpp deleted file mode 100644 index c841e612e..000000000 --- a/test_conformance/vulkan/test_vulkan_interop_buffer_fence.cpp +++ /dev/null @@ -1,1798 +0,0 @@ -// -// Copyright (c) 2022 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 -#include -#include -#include -#include -#include -#include "harness/errorHelpers.h" - -#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 { -struct Params -{ - uint32_t numBuffers; - uint32_t bufferSize; - uint32_t interBufferOffset; -}; -} - -static 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\ -}"; - -static 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\ -}"; - -static 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\ -}"; - - -static 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\ -}"; - -int run_test_with_two_queue_fence(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) -{ - int err = CL_SUCCESS; - size_t global_work_size[1]; - uint8_t *error_2; - cl_mem error_1; - cl_kernel update_buffer_kernel; - cl_kernel kernel_cq; - const char *program_source_const = kernel_text_numbuffer_2; - size_t program_source_length = strlen(program_source_const); - cl_program program = clCreateProgramWithSource( - context, 1, &program_source_const, &program_source_length, &err); - err = clBuildProgram(program, 0, NULL, NULL, NULL, NULL); - if (err != CL_SUCCESS) - { - print_error(err, "Error: Failed to build program \n"); - return err; - } - // create the kernel - kernel_cq = clCreateKernel(program, "clUpdateBuffer", &err); - if (err != CL_SUCCESS) - { - print_error(err, "clCreateKernel failed \n"); - return err; - } - - const std::vector - vkExternalMemoryHandleTypeList = - getSupportedVulkanExternalMemoryHandleTypeList(); - - VulkanQueue &vkQueue = vkDevice.getQueue(); - - std::vector vkBufferShader = readFile("buffer.spv"); - - VulkanShaderModule vkBufferShaderModule(vkDevice, vkBufferShader); - VulkanDescriptorSetLayoutBindingList vkDescriptorSetLayoutBindingList( - MAX_BUFFERS + 1, VULKAN_DESCRIPTOR_TYPE_STORAGE_BUFFER); - VulkanDescriptorSetLayout vkDescriptorSetLayout( - vkDevice, vkDescriptorSetLayoutBindingList); - VulkanPipelineLayout vkPipelineLayout(vkDevice, vkDescriptorSetLayout); - VulkanComputePipeline vkComputePipeline(vkDevice, vkPipelineLayout, - vkBufferShaderModule); - - VulkanDescriptorPool vkDescriptorPool(vkDevice, - vkDescriptorSetLayoutBindingList); - VulkanDescriptorSet vkDescriptorSet(vkDevice, vkDescriptorPool, - vkDescriptorSetLayout); - - const uint32_t maxIter = innerIterations; - VulkanCommandPool vkCommandPool(vkDevice); - VulkanCommandBuffer vkCommandBuffer(vkDevice, vkCommandPool); - - VulkanBuffer vkParamsBuffer(vkDevice, sizeof(Params)); - VulkanDeviceMemory vkParamsDeviceMemory( - vkDevice, vkParamsBuffer.getSize(), - getVulkanMemoryType(vkDevice, - VULKAN_MEMORY_TYPE_PROPERTY_HOST_VISIBLE_COHERENT)); - vkParamsDeviceMemory.bindBuffer(vkParamsBuffer); - std::vector vkBufferListDeviceMemory; - std::vector externalMemory; - - VkFence fence; - VkFenceCreateInfo fenceInfo{}; - fenceInfo.sType = VK_STRUCTURE_TYPE_FENCE_CREATE_INFO; - fenceInfo.pNext = nullptr; - fenceInfo.flags = 0; - - VkResult vkStatus = vkCreateFence(vkDevice, &fenceInfo, nullptr, &fence); - - if (vkStatus != VK_SUCCESS) - { - print_error(vkStatus, "Error: Failed create fence.\n"); - goto CLEANUP; - } - - for (size_t emhtIdx = 0; emhtIdx < vkExternalMemoryHandleTypeList.size(); - emhtIdx++) - { - VulkanExternalMemoryHandleType vkExternalMemoryHandleType = - vkExternalMemoryHandleTypeList[emhtIdx]; - log_info("External memory handle type: %d\n", - vkExternalMemoryHandleType); - - VulkanBuffer vkDummyBuffer(vkDevice, 4 * 1024, - vkExternalMemoryHandleType); - const VulkanMemoryTypeList &memoryTypeList = - vkDummyBuffer.getMemoryTypeList(); - - for (size_t mtIdx = 0; mtIdx < memoryTypeList.size(); mtIdx++) - { - const VulkanMemoryType &memoryType = memoryTypeList[mtIdx]; - - log_info("Memory type index: %d\n", (uint32_t)memoryType); - log_info("Memory type property: %d\n", - memoryType.getMemoryTypeProperty()); - - VulkanBufferList vkBufferList(numBuffers, vkDevice, bufferSize, - vkExternalMemoryHandleType); - - for (size_t bIdx = 0; bIdx < numBuffers; bIdx++) - { - vkBufferListDeviceMemory.push_back( - new VulkanDeviceMemory(vkDevice, bufferSize, memoryType, - vkExternalMemoryHandleType)); - externalMemory.push_back(new clExternalMemory( - vkBufferListDeviceMemory[bIdx], vkExternalMemoryHandleType, - 0, bufferSize, context, deviceId)); - } - cl_mem buffers[MAX_BUFFERS]; - 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); - buffers[bIdx] = externalMemory[bIdx]->getExternalMemoryBuffer(); - vkDescriptorSet.update((uint32_t)bIdx + 1, vkBufferList[bIdx]); - } - vkCommandBuffer.begin(); - vkCommandBuffer.bindPipeline(vkComputePipeline); - vkCommandBuffer.bindDescriptorSets( - vkComputePipeline, vkPipelineLayout, vkDescriptorSet); - vkCommandBuffer.dispatch(512, 1, 1); - vkCommandBuffer.end(); - - if (vkBufferList.size() == 2) - { - update_buffer_kernel = kernel[0]; - } - else if (vkBufferList.size() == 3) - { - update_buffer_kernel = kernel[1]; - } - else if (vkBufferList.size() == 5) - { - update_buffer_kernel = 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++) - { - - if (iter == 0) - { - vkQueue.submit(vkCommandBuffer, fence); - } - else - { - vkQueue.submit(vkCommandBuffer, fence); - } - vkWaitForFences(vkDevice, 1, &fence, VK_TRUE, UINT64_MAX); - - err = clSetKernelArg(update_buffer_kernel, 0, sizeof(uint32_t), - (void *)&bufferSize); - err |= clSetKernelArg(kernel_cq, 0, sizeof(uint32_t), - (void *)&bufferSize); - err |= clSetKernelArg(kernel_cq, 1, sizeof(cl_mem), - (void *)&(buffers[0])); - - for (int i = 0; i < vkBufferList.size() - 1; i++) - { - err |= - clSetKernelArg(update_buffer_kernel, i + 1, - sizeof(cl_mem), (void *)&(buffers[i])); - } - - err |= - clSetKernelArg(kernel_cq, 2, sizeof(cl_mem), - (void *)&(buffers[vkBufferList.size() - 1])); - - if (err != CL_SUCCESS) - { - print_error(err, - "Error: Failed to set arg values for kernel\n"); - goto CLEANUP; - } - cl_event first_launch; - - err = clEnqueueNDRangeKernel(cmd_queue1, update_buffer_kernel, - 1, NULL, global_work_size, NULL, 0, - NULL, &first_launch); - if (err != CL_SUCCESS) - { - print_error(err, - "Error: Failed to launch update_buffer_kernel," - "error\n"); - goto CLEANUP; - } - - err = clEnqueueNDRangeKernel(cmd_queue2, kernel_cq, 1, NULL, - global_work_size, NULL, 1, - &first_launch, NULL); - if (err != CL_SUCCESS) - { - print_error(err, - "Error: Failed to launch update_buffer_kernel," - "error\n"); - goto CLEANUP; - } - - if (iter != (maxIter - 1)) - { - vkWaitForFences(vkDevice, 1, &fence, VK_TRUE, UINT64_MAX); - } - } - error_2 = (uint8_t *)malloc(sizeof(uint8_t)); - if (NULL == error_2) - { - log_error("Not able to allocate memory\n"); - goto CLEANUP; - } - clFinish(cmd_queue2); - error_1 = clCreateBuffer(context, CL_MEM_WRITE_ONLY, - sizeof(uint8_t), NULL, &err); - if (CL_SUCCESS != err) - { - print_error(err, "Error: clCreateBuffer \n"); - goto CLEANUP; - } - uint8_t val = 0; - err = clEnqueueWriteBuffer(cmd_queue1, error_1, CL_TRUE, 0, - sizeof(uint8_t), &val, 0, NULL, NULL); - if (err != CL_SUCCESS) - { - print_error(err, "Error: Failed read output, error\n"); - goto CLEANUP; - } - - int calc_max_iter; - for (int i = 0; i < vkBufferList.size(); i++) - { - if (i == 0) - calc_max_iter = (maxIter * 3); - else - calc_max_iter = (maxIter * 2); - err = clSetKernelArg(verify_kernel, 0, sizeof(cl_mem), - (void *)&(buffers[i])); - 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); - if (err != CL_SUCCESS) - { - print_error(err, - "Error: Failed to set arg values for " - "verify_kernel \n"); - goto CLEANUP; - } - err = clEnqueueNDRangeKernel(cmd_queue1, verify_kernel, 1, NULL, - global_work_size, NULL, 0, NULL, - NULL); - - if (err != CL_SUCCESS) - { - print_error(err, - "Error: Failed to launch verify_kernel," - "error \n"); - goto CLEANUP; - } - err = clEnqueueReadBuffer(cmd_queue1, error_1, CL_TRUE, 0, - sizeof(uint8_t), error_2, 0, NULL, - NULL); - if (err != CL_SUCCESS) - { - print_error(err, "Error: Failed read output, error \n "); - goto CLEANUP; - } - if (*error_2 == 1) - { - log_error("&&&& vulkan_opencl_buffer test FAILED\n"); - goto CLEANUP; - } - } - for (size_t i = 0; i < vkBufferList.size(); i++) - { - delete vkBufferListDeviceMemory[i]; - delete externalMemory[i]; - } - vkBufferListDeviceMemory.erase(vkBufferListDeviceMemory.begin(), - vkBufferListDeviceMemory.begin() - + numBuffers); - externalMemory.erase(externalMemory.begin(), - externalMemory.begin() + numBuffers); - } - } -CLEANUP: - for (size_t i = 0; i < vkBufferListDeviceMemory.size(); i++) - { - if (vkBufferListDeviceMemory[i]) - { - delete vkBufferListDeviceMemory[i]; - } - if (externalMemory[i]) - { - delete externalMemory[i]; - } - } - if (program) clReleaseProgram(program); - if (kernel_cq) clReleaseKernel(kernel_cq); - if (error_2) free(error_2); - if (error_1) clReleaseMemObject(error_1); - - vkDestroyFence(vkDevice, fence, nullptr); - - return err; -} - -int run_test_with_one_queue_fence(cl_context &context, - cl_command_queue &cmd_queue1, - cl_kernel *kernel, cl_kernel &verify_kernel, - VulkanDevice &vkDevice, uint32_t numBuffers, - uint32_t bufferSize) -{ - log_info("RUNNING TEST WITH ONE QUEUE...... \n\n"); - size_t global_work_size[1]; - uint8_t *error_2; - cl_mem error_1; - cl_kernel update_buffer_kernel; - int err = CL_SUCCESS; - - const std::vector - vkExternalMemoryHandleTypeList = - getSupportedVulkanExternalMemoryHandleTypeList(); - - VulkanQueue &vkQueue = vkDevice.getQueue(); - - std::vector vkBufferShader = readFile("buffer.spv"); - VulkanShaderModule vkBufferShaderModule(vkDevice, vkBufferShader); - VulkanDescriptorSetLayoutBindingList vkDescriptorSetLayoutBindingList( - MAX_BUFFERS + 1, VULKAN_DESCRIPTOR_TYPE_STORAGE_BUFFER); - VulkanDescriptorSetLayout vkDescriptorSetLayout( - vkDevice, vkDescriptorSetLayoutBindingList); - VulkanPipelineLayout vkPipelineLayout(vkDevice, vkDescriptorSetLayout); - VulkanComputePipeline vkComputePipeline(vkDevice, vkPipelineLayout, - vkBufferShaderModule); - - VulkanDescriptorPool vkDescriptorPool(vkDevice, - vkDescriptorSetLayoutBindingList); - VulkanDescriptorSet vkDescriptorSet(vkDevice, vkDescriptorPool, - vkDescriptorSetLayout); - - const uint32_t maxIter = innerIterations; - VulkanCommandPool vkCommandPool(vkDevice); - VulkanCommandBuffer vkCommandBuffer(vkDevice, vkCommandPool); - - VulkanBuffer vkParamsBuffer(vkDevice, sizeof(Params)); - VulkanDeviceMemory vkParamsDeviceMemory( - vkDevice, vkParamsBuffer.getSize(), - getVulkanMemoryType(vkDevice, - VULKAN_MEMORY_TYPE_PROPERTY_HOST_VISIBLE_COHERENT)); - vkParamsDeviceMemory.bindBuffer(vkParamsBuffer); - std::vector vkBufferListDeviceMemory; - std::vector externalMemory; - - VkFence fence; - VkFenceCreateInfo fenceInfo{}; - fenceInfo.sType = VK_STRUCTURE_TYPE_FENCE_CREATE_INFO; - fenceInfo.pNext = nullptr; - fenceInfo.flags = 0; - - VkResult vkStatus = vkCreateFence(vkDevice, &fenceInfo, nullptr, &fence); - - if (vkStatus != VK_SUCCESS) - { - print_error(vkStatus, "Error: Failed create fence.\n"); - goto CLEANUP; - } - - for (size_t emhtIdx = 0; emhtIdx < vkExternalMemoryHandleTypeList.size(); - emhtIdx++) - { - VulkanExternalMemoryHandleType vkExternalMemoryHandleType = - vkExternalMemoryHandleTypeList[emhtIdx]; - log_info("External memory handle type: %d\n", - vkExternalMemoryHandleType); - - VulkanBuffer vkDummyBuffer(vkDevice, 4 * 1024, - vkExternalMemoryHandleType); - const VulkanMemoryTypeList &memoryTypeList = - vkDummyBuffer.getMemoryTypeList(); - - for (size_t mtIdx = 0; mtIdx < memoryTypeList.size(); mtIdx++) - { - const VulkanMemoryType &memoryType = memoryTypeList[mtIdx]; - - log_info("Memory type index: %d\n", (uint32_t)memoryType); - log_info("Memory type property: %d\n", - memoryType.getMemoryTypeProperty()); - - VulkanBufferList vkBufferList(numBuffers, vkDevice, bufferSize, - vkExternalMemoryHandleType); - - for (size_t bIdx = 0; bIdx < numBuffers; bIdx++) - { - vkBufferListDeviceMemory.push_back( - new VulkanDeviceMemory(vkDevice, bufferSize, memoryType, - vkExternalMemoryHandleType)); - externalMemory.push_back(new clExternalMemory( - vkBufferListDeviceMemory[bIdx], vkExternalMemoryHandleType, - 0, bufferSize, context, deviceId)); - } - cl_mem buffers[4]; - 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); - buffers[bIdx] = externalMemory[bIdx]->getExternalMemoryBuffer(); - vkDescriptorSet.update((uint32_t)bIdx + 1, vkBufferList[bIdx]); - } - vkCommandBuffer.begin(); - vkCommandBuffer.bindPipeline(vkComputePipeline); - vkCommandBuffer.bindDescriptorSets( - vkComputePipeline, vkPipelineLayout, vkDescriptorSet); - vkCommandBuffer.dispatch(512, 1, 1); - vkCommandBuffer.end(); - - if (vkBufferList.size() == 1) - { - update_buffer_kernel = kernel[0]; - } - else if (vkBufferList.size() == 2) - { - update_buffer_kernel = kernel[1]; - } - else if (vkBufferList.size() == 4) - { - update_buffer_kernel = 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++) - { - if (iter == 0) - { - vkQueue.submit(vkCommandBuffer, fence); - } - else - { - vkQueue.submit(vkCommandBuffer, fence); - } - vkWaitForFences(vkDevice, 1, &fence, VK_TRUE, UINT64_MAX); - - err = clSetKernelArg(update_buffer_kernel, 0, sizeof(uint32_t), - (void *)&bufferSize); - for (int i = 0; i < vkBufferList.size(); i++) - { - err |= - clSetKernelArg(update_buffer_kernel, i + 1, - sizeof(cl_mem), (void *)&(buffers[i])); - } - - if (err != CL_SUCCESS) - { - print_error(err, - "Error: Failed to set arg values for kernel\n"); - goto CLEANUP; - } - err = clEnqueueNDRangeKernel(cmd_queue1, update_buffer_kernel, - 1, NULL, global_work_size, NULL, 0, - NULL, NULL); - if (err != CL_SUCCESS) - { - print_error(err, - "Error: Failed to launch update_buffer_kernel," - " error\n"); - goto CLEANUP; - } - if (iter != (maxIter - 1)) - { - vkWaitForFences(vkDevice, 1, &fence, VK_TRUE, UINT64_MAX); - clFinish(cmd_queue1); - } - } - error_2 = (uint8_t *)malloc(sizeof(uint8_t)); - if (NULL == error_2) - { - log_error("Not able to allocate memory\n"); - goto CLEANUP; - } - - error_1 = clCreateBuffer(context, CL_MEM_WRITE_ONLY, - sizeof(uint8_t), NULL, &err); - if (CL_SUCCESS != err) - { - print_error(err, "Error: clCreateBuffer \n"); - goto CLEANUP; - } - uint8_t val = 0; - err = clEnqueueWriteBuffer(cmd_queue1, error_1, CL_TRUE, 0, - sizeof(uint8_t), &val, 0, NULL, NULL); - if (CL_SUCCESS != err) - { - print_error(err, "Error: clEnqueueWriteBuffer \n"); - goto CLEANUP; - } - - int calc_max_iter = (maxIter * 2); - for (int i = 0; i < vkBufferList.size(); i++) - { - err = clSetKernelArg(verify_kernel, 0, sizeof(cl_mem), - (void *)&(buffers[i])); - 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); - if (err != CL_SUCCESS) - { - print_error( - err, - "Error: Failed to set arg values for verify_kernel \n"); - goto CLEANUP; - } - err = clEnqueueNDRangeKernel(cmd_queue1, verify_kernel, 1, NULL, - global_work_size, NULL, 0, NULL, - NULL); - if (err != CL_SUCCESS) - { - print_error( - err, "Error: Failed to launch verify_kernel, error\n"); - goto CLEANUP; - } - - err = clEnqueueReadBuffer(cmd_queue1, error_1, CL_TRUE, 0, - sizeof(uint8_t), error_2, 0, NULL, - NULL); - if (err != CL_SUCCESS) - { - print_error(err, "Error: Failed read output, error \n"); - goto CLEANUP; - } - if (*error_2 == 1) - { - log_error("&&&& vulkan_opencl_buffer test FAILED\n"); - goto CLEANUP; - } - } - for (size_t i = 0; i < vkBufferList.size(); i++) - { - delete vkBufferListDeviceMemory[i]; - delete externalMemory[i]; - } - vkBufferListDeviceMemory.erase(vkBufferListDeviceMemory.begin(), - vkBufferListDeviceMemory.begin() - + numBuffers); - externalMemory.erase(externalMemory.begin(), - externalMemory.begin() + numBuffers); - } - } -CLEANUP: - for (size_t i = 0; i < vkBufferListDeviceMemory.size(); i++) - { - if (vkBufferListDeviceMemory[i]) - { - delete vkBufferListDeviceMemory[i]; - } - if (externalMemory[i]) - { - delete externalMemory[i]; - } - } - if (error_2) free(error_2); - if (error_1) clReleaseMemObject(error_1); - - vkDestroyFence(vkDevice, fence, nullptr); - - return err; -} - -int run_test_with_multi_import_same_ctx_fence( - 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) -{ - size_t global_work_size[1]; - uint8_t *error_2; - cl_mem error_1; - int numImports = numBuffers; - cl_kernel update_buffer_kernel[MAX_IMPORTS]; - int err = CL_SUCCESS; - int calc_max_iter; - bool withOffset; - uint32_t pBufferSize; - - const std::vector - vkExternalMemoryHandleTypeList = - getSupportedVulkanExternalMemoryHandleTypeList(); - - VulkanQueue &vkQueue = vkDevice.getQueue(); - - std::vector vkBufferShader = readFile("buffer.spv"); - - VulkanShaderModule vkBufferShaderModule(vkDevice, vkBufferShader); - VulkanDescriptorSetLayoutBindingList vkDescriptorSetLayoutBindingList( - MAX_BUFFERS + 1, VULKAN_DESCRIPTOR_TYPE_STORAGE_BUFFER); - VulkanDescriptorSetLayout vkDescriptorSetLayout( - vkDevice, vkDescriptorSetLayoutBindingList); - VulkanPipelineLayout vkPipelineLayout(vkDevice, vkDescriptorSetLayout); - VulkanComputePipeline vkComputePipeline(vkDevice, vkPipelineLayout, - vkBufferShaderModule); - - VulkanDescriptorPool vkDescriptorPool(vkDevice, - vkDescriptorSetLayoutBindingList); - VulkanDescriptorSet vkDescriptorSet(vkDevice, vkDescriptorPool, - vkDescriptorSetLayout); - - const uint32_t maxIter = innerIterations; - VulkanCommandPool vkCommandPool(vkDevice); - VulkanCommandBuffer vkCommandBuffer(vkDevice, vkCommandPool); - - VulkanBuffer vkParamsBuffer(vkDevice, sizeof(Params)); - VulkanDeviceMemory vkParamsDeviceMemory( - vkDevice, vkParamsBuffer.getSize(), - getVulkanMemoryType(vkDevice, - VULKAN_MEMORY_TYPE_PROPERTY_HOST_VISIBLE_COHERENT)); - vkParamsDeviceMemory.bindBuffer(vkParamsBuffer); - std::vector vkBufferListDeviceMemory; - std::vector> externalMemory; - - VkFence fence; - VkFenceCreateInfo fenceInfo{}; - fenceInfo.sType = VK_STRUCTURE_TYPE_FENCE_CREATE_INFO; - fenceInfo.pNext = nullptr; - fenceInfo.flags = 0; - - VkResult vkStatus = vkCreateFence(vkDevice, &fenceInfo, nullptr, &fence); - - if (vkStatus != VK_SUCCESS) - { - print_error(vkStatus, "Error: Failed create fence.\n"); - goto CLEANUP; - } - - for (size_t emhtIdx = 0; emhtIdx < vkExternalMemoryHandleTypeList.size(); - emhtIdx++) - { - VulkanExternalMemoryHandleType vkExternalMemoryHandleType = - vkExternalMemoryHandleTypeList[emhtIdx]; - log_info("External memory handle type: %d\n", - vkExternalMemoryHandleType); - - VulkanBuffer vkDummyBuffer(vkDevice, 4 * 1024, - vkExternalMemoryHandleType); - const VulkanMemoryTypeList &memoryTypeList = - vkDummyBuffer.getMemoryTypeList(); - - for (size_t mtIdx = 0; mtIdx < memoryTypeList.size(); mtIdx++) - { - const VulkanMemoryType &memoryType = memoryTypeList[mtIdx]; - - log_info("Memory type index: %d\n", (uint32_t)memoryType); - log_info("Memory type property: %d\n", - memoryType.getMemoryTypeProperty()); - for (unsigned int withOffset = 0; - withOffset <= (unsigned int)enableOffset; withOffset++) - { - log_info("Running withOffset case %d\n", (uint32_t)withOffset); - if (withOffset) - { - pBufferSize = bufferSizeForOffset; - } - else - { - pBufferSize = bufferSize; - } - cl_mem buffers[MAX_BUFFERS][MAX_IMPORTS]; - VulkanBufferList vkBufferList(numBuffers, vkDevice, pBufferSize, - vkExternalMemoryHandleType); - uint32_t interBufferOffset = - (uint32_t)(vkBufferList[0].getSize()); - - for (size_t bIdx = 0; bIdx < numBuffers; bIdx++) - { - if (withOffset == 0) - { - vkBufferListDeviceMemory.push_back( - new VulkanDeviceMemory(vkDevice, pBufferSize, - memoryType, - vkExternalMemoryHandleType)); - } - if (withOffset == 1) - { - uint32_t totalSize = - (uint32_t)(vkBufferList.size() * interBufferOffset); - vkBufferListDeviceMemory.push_back( - new VulkanDeviceMemory(vkDevice, totalSize, - memoryType, - vkExternalMemoryHandleType)); - } - std::vector pExternalMemory; - for (size_t cl_bIdx = 0; cl_bIdx < numImports; cl_bIdx++) - { - pExternalMemory.push_back(new clExternalMemory( - vkBufferListDeviceMemory[bIdx], - vkExternalMemoryHandleType, - withOffset * bIdx * interBufferOffset, pBufferSize, - context, deviceId)); - } - externalMemory.push_back(pExternalMemory); - } - - clFinish(cmd_queue1); - Params *params = (Params *)vkParamsDeviceMemory.map(); - params->numBuffers = numBuffers; - params->bufferSize = pBufferSize; - params->interBufferOffset = interBufferOffset * withOffset; - vkParamsDeviceMemory.unmap(); - vkDescriptorSet.update(0, vkParamsBuffer); - for (size_t bIdx = 0; bIdx < vkBufferList.size(); bIdx++) - { - size_t buffer_size = vkBufferList[bIdx].getSize(); - vkBufferListDeviceMemory[bIdx]->bindBuffer( - vkBufferList[bIdx], - bIdx * interBufferOffset * withOffset); - for (size_t cl_bIdx = 0; cl_bIdx < numImports; cl_bIdx++) - { - buffers[bIdx][cl_bIdx] = - externalMemory[bIdx][cl_bIdx] - ->getExternalMemoryBuffer(); - } - vkDescriptorSet.update((uint32_t)bIdx + 1, - vkBufferList[bIdx]); - } - vkCommandBuffer.begin(); - vkCommandBuffer.bindPipeline(vkComputePipeline); - vkCommandBuffer.bindDescriptorSets( - vkComputePipeline, vkPipelineLayout, vkDescriptorSet); - vkCommandBuffer.dispatch(512, 1, 1); - vkCommandBuffer.end(); - for (int i = 0; i < numImports; i++) - { - update_buffer_kernel[i] = (numBuffers == 1) - ? kernel[0] - : ((numBuffers == 2) ? kernel[1] : kernel[2]); - } - // global work size should be less than or equal to - // bufferSizeList[i] - global_work_size[0] = pBufferSize; - - for (uint32_t iter = 0; iter < maxIter; iter++) - { - if (iter == 0) - { - vkQueue.submit(vkCommandBuffer, fence); - } - else - { - vkQueue.submit(vkCommandBuffer, fence); - } - vkWaitForFences(vkDevice, 1, &fence, VK_TRUE, UINT64_MAX); - for (uint8_t launchIter = 0; launchIter < numImports; - launchIter++) - { - err = clSetKernelArg(update_buffer_kernel[launchIter], - 0, sizeof(uint32_t), - (void *)&pBufferSize); - for (int i = 0; i < numBuffers; i++) - { - err |= clSetKernelArg( - update_buffer_kernel[launchIter], i + 1, - sizeof(cl_mem), - (void *)&(buffers[i][launchIter])); - } - - if (err != CL_SUCCESS) - { - print_error(err, - "Error: Failed to set arg values for " - "kernel\n "); - goto CLEANUP; - } - err = clEnqueueNDRangeKernel( - cmd_queue1, update_buffer_kernel[launchIter], 1, - NULL, global_work_size, NULL, 0, NULL, NULL); - if (err != CL_SUCCESS) - { - print_error(err, - "Error: Failed to launch " - "update_buffer_kernel, error\n "); - goto CLEANUP; - } - } - if (iter != (maxIter - 1)) - { - vkWaitForFences(vkDevice, 1, &fence, VK_TRUE, - UINT64_MAX); - } - } - error_2 = (uint8_t *)malloc(sizeof(uint8_t)); - if (NULL == error_2) - { - log_error("Not able to allocate memory\n"); - goto CLEANUP; - } - - error_1 = clCreateBuffer(context, CL_MEM_WRITE_ONLY, - sizeof(uint8_t), NULL, &err); - if (CL_SUCCESS != err) - { - print_error(err, "Error: clCreateBuffer \n"); - goto CLEANUP; - } - uint8_t val = 0; - err = - clEnqueueWriteBuffer(cmd_queue1, error_1, CL_TRUE, 0, - sizeof(uint8_t), &val, 0, NULL, NULL); - if (CL_SUCCESS != err) - { - print_error(err, "Error: clEnqueueWriteBuffer \n"); - goto CLEANUP; - } - calc_max_iter = maxIter * (numBuffers + 1); - - for (int i = 0; i < vkBufferList.size(); i++) - { - err = clSetKernelArg(verify_kernel, 0, sizeof(cl_mem), - (void *)&(buffers[i][0])); - err |= clSetKernelArg(verify_kernel, 1, sizeof(int), - &pBufferSize); - err |= clSetKernelArg(verify_kernel, 2, sizeof(int), - &calc_max_iter); - err |= clSetKernelArg(verify_kernel, 3, sizeof(cl_mem), - (void *)&error_1); - if (err != CL_SUCCESS) - { - print_error(err, - "Error: Failed to set arg values for " - "verify_kernel \n"); - goto CLEANUP; - } - err = clEnqueueNDRangeKernel(cmd_queue1, verify_kernel, 1, - NULL, global_work_size, NULL, - 0, NULL, NULL); - if (err != CL_SUCCESS) - { - print_error( - err, - "Error: Failed to launch verify_kernel, error\n"); - goto CLEANUP; - } - - err = clEnqueueReadBuffer(cmd_queue1, error_1, CL_TRUE, 0, - sizeof(uint8_t), error_2, 0, NULL, - NULL); - if (err != CL_SUCCESS) - { - print_error(err, "Error: Failed read output, error \n"); - goto CLEANUP; - } - if (*error_2 == 1) - { - log_error("&&&& vulkan_opencl_buffer test FAILED\n"); - goto CLEANUP; - } - } - 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: - for (size_t i = 0; i < vkBufferListDeviceMemory.size(); i++) - { - if (vkBufferListDeviceMemory[i]) - { - delete vkBufferListDeviceMemory[i]; - } - } - for (size_t i = 0; i < externalMemory.size(); i++) - { - for (size_t j = 0; j < externalMemory[i].size(); j++) - { - if (externalMemory[i][j]) - { - delete externalMemory[i][j]; - } - } - } - if (error_2) free(error_2); - if (error_1) clReleaseMemObject(error_1); - - vkDestroyFence(vkDevice, fence, nullptr); - - return err; -} - -int run_test_with_multi_import_diff_ctx_fence( - 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) -{ - size_t global_work_size[1]; - uint8_t *error_3; - cl_mem error_1; - cl_mem error_2; - int numImports = numBuffers; - cl_kernel update_buffer_kernel1[MAX_IMPORTS]; - cl_kernel update_buffer_kernel2[MAX_IMPORTS]; - - int err = CL_SUCCESS; - int calc_max_iter; - bool withOffset; - uint32_t pBufferSize; - - const std::vector - vkExternalMemoryHandleTypeList = - getSupportedVulkanExternalMemoryHandleTypeList(); - VulkanExternalSemaphoreHandleType vkExternalSemaphoreHandleType = - getSupportedVulkanExternalSemaphoreHandleTypeList()[0]; - - VulkanQueue &vkQueue = vkDevice.getQueue(); - - std::vector vkBufferShader = readFile("buffer.spv"); - - VulkanShaderModule vkBufferShaderModule(vkDevice, vkBufferShader); - VulkanDescriptorSetLayoutBindingList vkDescriptorSetLayoutBindingList( - MAX_BUFFERS + 1, VULKAN_DESCRIPTOR_TYPE_STORAGE_BUFFER); - VulkanDescriptorSetLayout vkDescriptorSetLayout( - vkDevice, vkDescriptorSetLayoutBindingList); - VulkanPipelineLayout vkPipelineLayout(vkDevice, vkDescriptorSetLayout); - VulkanComputePipeline vkComputePipeline(vkDevice, vkPipelineLayout, - vkBufferShaderModule); - - VulkanDescriptorPool vkDescriptorPool(vkDevice, - vkDescriptorSetLayoutBindingList); - VulkanDescriptorSet vkDescriptorSet(vkDevice, vkDescriptorPool, - vkDescriptorSetLayout); - - const uint32_t maxIter = innerIterations; - VulkanCommandPool vkCommandPool(vkDevice); - VulkanCommandBuffer vkCommandBuffer(vkDevice, vkCommandPool); - - VulkanBuffer vkParamsBuffer(vkDevice, sizeof(Params)); - VulkanDeviceMemory vkParamsDeviceMemory( - vkDevice, vkParamsBuffer.getSize(), - getVulkanMemoryType(vkDevice, - VULKAN_MEMORY_TYPE_PROPERTY_HOST_VISIBLE_COHERENT)); - vkParamsDeviceMemory.bindBuffer(vkParamsBuffer); - std::vector vkBufferListDeviceMemory; - std::vector> externalMemory1; - std::vector> externalMemory2; - - VkFence fence; - VkFenceCreateInfo fenceInfo{}; - fenceInfo.sType = VK_STRUCTURE_TYPE_FENCE_CREATE_INFO; - fenceInfo.pNext = nullptr; - fenceInfo.flags = 0; - - VkResult vkStatus = vkCreateFence(vkDevice, &fenceInfo, nullptr, &fence); - if (vkStatus != VK_SUCCESS) - { - print_error(vkStatus, "Error: Failed create fence.\n"); - goto CLEANUP; - } - - for (size_t emhtIdx = 0; emhtIdx < vkExternalMemoryHandleTypeList.size(); - emhtIdx++) - { - VulkanExternalMemoryHandleType vkExternalMemoryHandleType = - vkExternalMemoryHandleTypeList[emhtIdx]; - log_info("External memory handle type:%d\n", - vkExternalMemoryHandleType); - - VulkanBuffer vkDummyBuffer(vkDevice, 4 * 1024, - vkExternalMemoryHandleType); - const VulkanMemoryTypeList &memoryTypeList = - vkDummyBuffer.getMemoryTypeList(); - - for (size_t mtIdx = 0; mtIdx < memoryTypeList.size(); mtIdx++) - { - const VulkanMemoryType &memoryType = memoryTypeList[mtIdx]; - - log_info("Memory type index: %d\n", (uint32_t)memoryType); - log_info("Memory type property: %d\n", - memoryType.getMemoryTypeProperty()); - - for (unsigned int withOffset = 0; - withOffset <= (unsigned int)enableOffset; withOffset++) - { - log_info("Running withOffset case %d\n", (uint32_t)withOffset); - cl_mem buffers1[MAX_BUFFERS][MAX_IMPORTS]; - cl_mem buffers2[MAX_BUFFERS][MAX_IMPORTS]; - if (withOffset) - { - pBufferSize = bufferSizeForOffset; - } - else - { - pBufferSize = bufferSize; - } - VulkanBufferList vkBufferList(numBuffers, vkDevice, pBufferSize, - vkExternalMemoryHandleType); - uint32_t interBufferOffset = - (uint32_t)(vkBufferList[0].getSize()); - - for (size_t bIdx = 0; bIdx < numBuffers; bIdx++) - { - if (withOffset == 0) - { - vkBufferListDeviceMemory.push_back( - new VulkanDeviceMemory(vkDevice, pBufferSize, - memoryType, - vkExternalMemoryHandleType)); - } - if (withOffset == 1) - { - uint32_t totalSize = - (uint32_t)(vkBufferList.size() * interBufferOffset); - vkBufferListDeviceMemory.push_back( - new VulkanDeviceMemory(vkDevice, totalSize, - memoryType, - vkExternalMemoryHandleType)); - } - std::vector pExternalMemory1; - std::vector pExternalMemory2; - for (size_t cl_bIdx = 0; cl_bIdx < numImports; cl_bIdx++) - { - pExternalMemory1.push_back(new clExternalMemory( - vkBufferListDeviceMemory[bIdx], - vkExternalMemoryHandleType, - withOffset * bIdx * interBufferOffset, pBufferSize, - context, deviceId)); - pExternalMemory2.push_back(new clExternalMemory( - vkBufferListDeviceMemory[bIdx], - vkExternalMemoryHandleType, - withOffset * bIdx * interBufferOffset, pBufferSize, - context2, deviceId)); - } - externalMemory1.push_back(pExternalMemory1); - externalMemory2.push_back(pExternalMemory2); - } - - clFinish(cmd_queue1); - Params *params = (Params *)vkParamsDeviceMemory.map(); - params->numBuffers = numBuffers; - params->bufferSize = pBufferSize; - params->interBufferOffset = interBufferOffset * withOffset; - vkParamsDeviceMemory.unmap(); - vkDescriptorSet.update(0, vkParamsBuffer); - for (size_t bIdx = 0; bIdx < vkBufferList.size(); bIdx++) - { - size_t buffer_size = vkBufferList[bIdx].getSize(); - vkBufferListDeviceMemory[bIdx]->bindBuffer( - vkBufferList[bIdx], - bIdx * interBufferOffset * withOffset); - for (size_t cl_bIdx = 0; cl_bIdx < numImports; cl_bIdx++) - { - buffers1[bIdx][cl_bIdx] = - externalMemory1[bIdx][cl_bIdx] - ->getExternalMemoryBuffer(); - buffers2[bIdx][cl_bIdx] = - externalMemory2[bIdx][cl_bIdx] - ->getExternalMemoryBuffer(); - } - vkDescriptorSet.update((uint32_t)bIdx + 1, - vkBufferList[bIdx]); - } - - vkCommandBuffer.begin(); - vkCommandBuffer.bindPipeline(vkComputePipeline); - vkCommandBuffer.bindDescriptorSets( - vkComputePipeline, vkPipelineLayout, vkDescriptorSet); - vkCommandBuffer.dispatch(512, 1, 1); - vkCommandBuffer.end(); - - for (int i = 0; i < numImports; i++) - { - update_buffer_kernel1[i] = (numBuffers == 1) - ? kernel1[0] - : ((numBuffers == 2) ? kernel1[1] : kernel1[2]); - update_buffer_kernel2[i] = (numBuffers == 1) - ? kernel2[0] - : ((numBuffers == 2) ? kernel2[1] : kernel2[2]); - } - - // global work size should be less than or equal - // to bufferSizeList[i] - global_work_size[0] = pBufferSize; - - for (uint32_t iter = 0; iter < maxIter; iter++) - { - if (iter == 0) - { - vkQueue.submit(vkCommandBuffer, fence); - } - else - { - vkQueue.submit(vkCommandBuffer, fence); - } - vkWaitForFences(vkDevice, 1, &fence, VK_TRUE, UINT64_MAX); - - for (uint8_t launchIter = 0; launchIter < numImports; - launchIter++) - { - err = clSetKernelArg(update_buffer_kernel1[launchIter], - 0, sizeof(uint32_t), - (void *)&pBufferSize); - for (int i = 0; i < numBuffers; i++) - { - err |= clSetKernelArg( - update_buffer_kernel1[launchIter], i + 1, - sizeof(cl_mem), - (void *)&(buffers1[i][launchIter])); - } - - if (err != CL_SUCCESS) - { - print_error(err, - "Error: Failed to set arg values for " - "kernel\n "); - goto CLEANUP; - } - err = clEnqueueNDRangeKernel( - cmd_queue1, update_buffer_kernel1[launchIter], 1, - NULL, global_work_size, NULL, 0, NULL, NULL); - if (err != CL_SUCCESS) - { - print_error(err, - "Error: Failed to launch " - "update_buffer_kernel, error\n"); - goto CLEANUP; - } - } - if (iter != (maxIter - 1)) - { - vkWaitForFences(vkDevice, 1, &fence, VK_TRUE, - UINT64_MAX); - } - } - clFinish(cmd_queue1); - for (uint32_t iter = 0; iter < maxIter; iter++) - { - if (iter == 0) - { - vkQueue.submit(vkCommandBuffer, fence); - } - else - { - vkQueue.submit(vkCommandBuffer, fence); - } - vkWaitForFences(vkDevice, 1, &fence, VK_TRUE, UINT64_MAX); - - for (uint8_t launchIter = 0; launchIter < numImports; - launchIter++) - { - err = clSetKernelArg(update_buffer_kernel2[launchIter], - 0, sizeof(uint32_t), - (void *)&bufferSize); - for (int i = 0; i < numBuffers; i++) - { - err |= clSetKernelArg( - update_buffer_kernel2[launchIter], i + 1, - sizeof(cl_mem), - (void *)&(buffers2[i][launchIter])); - } - - if (err != CL_SUCCESS) - { - print_error(err, - "Error: Failed to set arg values for " - "kernel\n "); - goto CLEANUP; - } - err = clEnqueueNDRangeKernel( - cmd_queue2, update_buffer_kernel2[launchIter], 1, - NULL, global_work_size, NULL, 0, NULL, NULL); - if (err != CL_SUCCESS) - { - print_error(err, - "Error: Failed to launch " - "update_buffer_kernel, error\n "); - goto CLEANUP; - } - } - if (iter != (maxIter - 1)) - { - vkWaitForFences(vkDevice, 1, &fence, VK_TRUE, - UINT64_MAX); - } - } - clFinish(cmd_queue2); - error_3 = (uint8_t *)malloc(sizeof(uint8_t)); - if (NULL == error_3) - { - log_error("Not able to allocate memory\n"); - goto CLEANUP; - } - - error_1 = clCreateBuffer(context, CL_MEM_WRITE_ONLY, - sizeof(uint8_t), NULL, &err); - if (CL_SUCCESS != err) - { - print_error(err, "Error: clCreateBuffer \n"); - goto CLEANUP; - } - error_2 = clCreateBuffer(context2, CL_MEM_WRITE_ONLY, - sizeof(uint8_t), NULL, &err); - if (CL_SUCCESS != err) - { - print_error(err, "Error: clCreateBuffer \n"); - goto CLEANUP; - } - uint8_t val = 0; - err = - clEnqueueWriteBuffer(cmd_queue1, error_1, CL_TRUE, 0, - sizeof(uint8_t), &val, 0, NULL, NULL); - if (err != CL_SUCCESS) - { - print_error(err, "Error: Failed read output, error \n"); - goto CLEANUP; - } - - err = - clEnqueueWriteBuffer(cmd_queue2, error_2, CL_TRUE, 0, - sizeof(uint8_t), &val, 0, NULL, NULL); - if (err != CL_SUCCESS) - { - print_error(err, "Error: Failed read output, error \n"); - goto CLEANUP; - } - - 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); - if (err != CL_SUCCESS) - { - print_error(err, - "Error: Failed to set arg values for " - "verify_kernel \n"); - goto CLEANUP; - } - err = clEnqueueNDRangeKernel(cmd_queue1, verify_kernel, 1, - NULL, global_work_size, NULL, - 0, NULL, NULL); - if (err != CL_SUCCESS) - { - print_error(err, - "Error: Failed to launch verify_kernel," - "error\n"); - goto CLEANUP; - } - - err = clEnqueueReadBuffer(cmd_queue1, error_1, CL_TRUE, 0, - sizeof(uint8_t), error_3, 0, NULL, - NULL); - if (err != CL_SUCCESS) - { - print_error(err, "Error: Failed read output, error\n"); - goto CLEANUP; - } - if (*error_3 == 1) - { - log_error("&&&& vulkan_opencl_buffer test FAILED\n"); - goto CLEANUP; - } - } - *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); - if (err != CL_SUCCESS) - { - print_error(err, - "Error: Failed to set arg values for " - "verify_kernel \n"); - goto CLEANUP; - } - err = clEnqueueNDRangeKernel(cmd_queue2, verify_kernel2, 1, - NULL, global_work_size, NULL, - 0, NULL, NULL); - if (err != CL_SUCCESS) - { - print_error(err, - "Error: Failed to launch verify_kernel," - "error\n"); - goto CLEANUP; - } - - err = clEnqueueReadBuffer(cmd_queue2, error_2, CL_TRUE, 0, - sizeof(uint8_t), error_3, 0, NULL, - NULL); - if (err != CL_SUCCESS) - { - print_error(err, "Error: Failed read output, error\n"); - goto CLEANUP; - } - if (*error_3 == 1) - { - log_error("&&&& vulkan_opencl_buffer test FAILED\n"); - goto CLEANUP; - } - } - 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(); - } - } - } -CLEANUP: - for (size_t i = 0; i < vkBufferListDeviceMemory.size(); i++) - { - if (vkBufferListDeviceMemory[i]) - { - delete vkBufferListDeviceMemory[i]; - } - } - for (size_t i = 0; i < externalMemory1.size(); i++) - { - for (size_t j = 0; j < externalMemory1[i].size(); j++) - { - if (externalMemory1[i][j]) - { - delete externalMemory1[i][j]; - } - } - } - for (size_t i = 0; i < externalMemory2.size(); i++) - { - for (size_t j = 0; j < externalMemory2[i].size(); j++) - { - if (externalMemory2[i][j]) - { - delete externalMemory2[i][j]; - } - } - } - if (error_3) free(error_3); - if (error_1) clReleaseMemObject(error_1); - if (error_2) clReleaseMemObject(error_2); - - vkDestroyFence(vkDevice, fence, nullptr); - - return err; -} - -int test_buffer_common_fence(cl_device_id device_, cl_context context_, - cl_command_queue queue_, int numElements_) -{ - - 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 }; - errNum = clGetPlatformIDs(1, &platform, NULL); - if (errNum != CL_SUCCESS) - { - print_error(errNum, "Error: Failed to get platform\n"); - goto CLEANUP; - } - - errNum = - clGetDeviceIDs(platform, CL_DEVICE_TYPE_GPU, 0, NULL, &num_devices); - if (CL_SUCCESS != errNum) - { - print_error(errNum, "clGetDeviceIDs failed in returning of devices\n"); - goto CLEANUP; - } - devices = (cl_device_id *)malloc(num_devices * sizeof(cl_device_id)); - if (NULL == devices) - { - errNum = CL_OUT_OF_HOST_MEMORY; - print_error(errNum, "Unable to allocate memory for devices\n"); - goto CLEANUP; - } - errNum = clGetDeviceIDs(platform, CL_DEVICE_TYPE_GPU, num_devices, devices, - NULL); - if (CL_SUCCESS != errNum) - { - print_error(errNum, "Failed to get deviceID.\n"); - goto CLEANUP; - } - contextProperties[1] = (cl_context_properties)platform; - log_info("Assigned contextproperties for platform\n"); - for (device_no = 0; device_no < num_devices; device_no++) - { - errNum = clGetDeviceInfo(devices[device_no], CL_DEVICE_EXTENSIONS, 0, - NULL, &extensionSize); - if (CL_SUCCESS != errNum) - { - print_error(errNum, - "Error in clGetDeviceInfo for getting device_extension " - "size....\n"); - goto CLEANUP; - } - extensions = (char *)malloc(extensionSize); - if (NULL == extensions) - { - print_error(errNum, "Unable to allocate memory for extensions\n"); - errNum = CL_OUT_OF_HOST_MEMORY; - goto CLEANUP; - } - errNum = clGetDeviceInfo(devices[device_no], CL_DEVICE_EXTENSIONS, - extensionSize, extensions, NULL); - if (CL_SUCCESS != errNum) - { - print_error(errNum, - "Error in clGetDeviceInfo for device_extension\n"); - goto CLEANUP; - } - errNum = clGetDeviceInfo(devices[device_no], CL_DEVICE_UUID_KHR, - CL_UUID_SIZE_KHR, uuid, &extensionSize); - if (CL_SUCCESS != errNum) - { - print_error(errNum, "clGetDeviceInfo failed\n"); - goto CLEANUP; - } - errNum = - memcmp(uuid, vkDevice.getPhysicalDevice().getUUID(), VK_UUID_SIZE); - if (errNum == 0) - { - break; - } - } - if (device_no >= num_devices) - { - errNum = EXIT_FAILURE; - print_error(errNum, - "OpenCL error: " - "No Vulkan-OpenCL Interop capable GPU found.\n"); - goto CLEANUP; - } - deviceId = devices[device_no]; - context = clCreateContextFromType(contextProperties, CL_DEVICE_TYPE_GPU, - NULL, NULL, &errNum); - if (CL_SUCCESS != errNum) - { - print_error(errNum, "error creating context\n"); - goto CLEANUP; - } - log_info("Successfully created context !!!\n"); - - cmd_queue1 = clCreateCommandQueue(context, devices[device_no], 0, &errNum); - if (CL_SUCCESS != errNum) - { - errNum = CL_INVALID_COMMAND_QUEUE; - print_error(errNum, "Error: Failed to create command queue!\n"); - goto CLEANUP; - } - cmd_queue2 = clCreateCommandQueue(context, devices[device_no], 0, &errNum); - if (CL_SUCCESS != errNum) - { - errNum = CL_INVALID_COMMAND_QUEUE; - print_error(errNum, "Error: Failed to create command queue!\n"); - goto CLEANUP; - } - 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); - if (errNum != CL_SUCCESS) - { - print_error(errNum, "Error: Failed to build program \n"); - return errNum; - } - // create the kernel - kernel[i] = clCreateKernel(program[i], "clUpdateBuffer", &errNum); - if (errNum != CL_SUCCESS) - { - print_error(errNum, "clCreateKernel failed \n"); - return errNum; - } - } - - 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); - if (errNum != CL_SUCCESS) - { - log_error("Error: Failed to build program2\n"); - return errNum; - } - verify_kernel = clCreateKernel(program_verify, "checkKernel", &errNum); - if (errNum != CL_SUCCESS) - { - print_error(errNum, "clCreateKernel failed \n"); - return errNum; - } - - if (multiCtx) // different context guard - { - context2 = clCreateContextFromType( - contextProperties, CL_DEVICE_TYPE_GPU, NULL, NULL, &errNum); - if (CL_SUCCESS != errNum) - { - print_error(errNum, "error creating context\n"); - goto CLEANUP; - } - cmd_queue3 = - clCreateCommandQueue(context2, devices[device_no], 0, &errNum); - if (CL_SUCCESS != errNum) - { - errNum = CL_INVALID_COMMAND_QUEUE; - print_error(errNum, "Error: Failed to create command queue!\n"); - goto CLEANUP; - } - for (int i = 0; i < 3; i++) - { - 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); - if (errNum != CL_SUCCESS) - { - print_error(errNum, "Error: Failed to build program \n"); - return errNum; - } - // create the kernel - kernel2[i] = clCreateKernel(program[i], "clUpdateBuffer", &errNum); - if (errNum != CL_SUCCESS) - { - print_error(errNum, "clCreateKernel failed \n"); - return errNum; - } - } - program_source_length = strlen(program_source_const_verify); - program_verify = - clCreateProgramWithSource(context2, 1, &program_source_const_verify, - &program_source_length, &errNum); - errNum = clBuildProgram(program_verify, 0, NULL, NULL, NULL, NULL); - if (errNum != CL_SUCCESS) - { - log_error("Error: Failed to build program2\n"); - return errNum; - } - verify_kernel2 = clCreateKernel(program_verify, "checkKernel", &errNum); - if (errNum != CL_SUCCESS) - { - print_error(errNum, "clCreateKernel failed \n"); - return errNum; - } - } - - 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++) - { - uint32_t bufferSize = bufferSizeList[sizeIdx]; - uint32_t bufferSizeForOffset = bufferSizeListforOffset[sizeIdx]; - log_info("&&&& RUNNING vulkan_opencl_buffer test for Buffer size: " - "%d\n", - bufferSize); - if (multiImport && !multiCtx) - { - errNum = run_test_with_multi_import_same_ctx_fence( - context, cmd_queue1, kernel, verify_kernel, vkDevice, - numBuffers, bufferSize, bufferSizeForOffset); - } - else if (multiImport && multiCtx) - { - errNum = run_test_with_multi_import_diff_ctx_fence( - context, context2, cmd_queue1, cmd_queue3, kernel, kernel2, - verify_kernel, verify_kernel2, vkDevice, numBuffers, - bufferSize, bufferSizeForOffset); - } - else if (numCQ == 2) - { - errNum = run_test_with_two_queue_fence( - context, cmd_queue1, cmd_queue2, kernel, verify_kernel, - vkDevice, numBuffers + 1, bufferSize); - } - else - { - errNum = run_test_with_one_queue_fence( - context, cmd_queue1, kernel, verify_kernel, vkDevice, - numBuffers, bufferSize); - } - if (errNum != CL_SUCCESS) - { - print_error(errNum, "func_name failed \n"); - goto CLEANUP; - } - } - } - -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); - - if (devices) free(devices); - if (extensions) free(extensions); - - return errNum; -} From f2466e47a1edc1abbde188b332a8ebfb64bff275 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Pawe=C5=82=20Jastrz=C4=99bski?= Date: Fri, 17 Mar 2023 17:26:51 +0100 Subject: [PATCH 04/14] Fix build break. MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Signed-off-by: Paweł Jastrzębski --- test_conformance/vulkan/main.cpp | 37 ++++++++++++++------------------ 1 file changed, 16 insertions(+), 21 deletions(-) diff --git a/test_conformance/vulkan/main.cpp b/test_conformance/vulkan/main.cpp index 46a96ecd1..c37d61637 100644 --- a/test_conformance/vulkan/main.cpp +++ b/test_conformance/vulkan/main.cpp @@ -58,8 +58,7 @@ 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_, - bool fence) + cl_command_queue queue_, int numElements_) { params_reset(); log_info("RUNNING TEST WITH ONE QUEUE...... \n\n"); @@ -146,25 +145,21 @@ int test_image_multiple_queue(cl_device_id device_, cl_context context_, 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), - 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), - ADD_TEST(consistency_external_image), - ADD_TEST(consistency_external_semaphore), - ADD_TEST(platform_info), - ADD_TEST(device_info) -}; +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), + ADD_TEST(consistency_external_image), + ADD_TEST(consistency_external_semaphore), + ADD_TEST(platform_info), + ADD_TEST(device_info) }; const int test_num = ARRAY_SIZE(test_list); From 92b6e1cbee769e75fc916287fbf8e3ef8ab6edbc Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Pawe=C5=82=20Jastrz=C4=99bski?= Date: Fri, 17 Mar 2023 17:37:03 +0100 Subject: [PATCH 05/14] Fix resource release conditions. MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Signed-off-by: Paweł Jastrzębski --- .../vulkan/test_vulkan_interop_buffer.cpp | 22 ++++--------------- 1 file changed, 4 insertions(+), 18 deletions(-) diff --git a/test_conformance/vulkan/test_vulkan_interop_buffer.cpp b/test_conformance/vulkan/test_vulkan_interop_buffer.cpp index 488a904b7..f623f8c0b 100644 --- a/test_conformance/vulkan/test_vulkan_interop_buffer.cpp +++ b/test_conformance/vulkan/test_vulkan_interop_buffer.cpp @@ -440,11 +440,7 @@ 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 (use_fence) - { - vkDestroyFence(vkDevice, fence, nullptr); - } - else + if (!use_fence) { if (clVk2CLExternalSemaphore) delete clVk2CLExternalSemaphore; if (clCl2VkExternalSemaphore) delete clCl2VkExternalSemaphore; @@ -770,11 +766,7 @@ int run_test_with_one_queue(cl_context &context, cl_command_queue &cmd_queue1, } } - if (use_fence) - { - vkDestroyFence(vkDevice, fence, nullptr); - } - else + if (!use_fence) { if (clVk2CLExternalSemaphore) delete clVk2CLExternalSemaphore; if (clCl2VkExternalSemaphore) delete clCl2VkExternalSemaphore; @@ -1164,11 +1156,7 @@ int run_test_with_multi_import_same_ctx( } } - if (use_fence) - { - vkDestroyFence(vkDevice, fence, nullptr); - } - else + if (!use_fence) { if (clVk2CLExternalSemaphore) delete clVk2CLExternalSemaphore; if (clCl2VkExternalSemaphore) delete clCl2VkExternalSemaphore; @@ -1745,9 +1733,7 @@ int run_test_with_multi_import_diff_ctx( } } - if (use_fence) - {} - else + if (!use_fence) { if (clVk2CLExternalSemaphore) delete clVk2CLExternalSemaphore; if (clCl2VkExternalSemaphore) delete clCl2VkExternalSemaphore; From 3e4f49b0caa11f124bd0b7e58404993b70c1f265 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Pawe=C5=82=20Jastrz=C4=99bski?= Date: Wed, 22 Mar 2023 17:03:10 +0100 Subject: [PATCH 06/14] Fix fence usage. MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Fixed following fence issues: - Add missing link to command buffer - Add fence reset before wait Signed-off-by: Paweł Jastrzębski --- .../common/vulkan_wrapper/vulkan_wrapper.cpp | 5 +++++ test_conformance/vulkan/test_vulkan_interop_buffer.cpp | 9 +++++++++ 2 files changed, 14 insertions(+) diff --git a/test_conformance/common/vulkan_wrapper/vulkan_wrapper.cpp b/test_conformance/common/vulkan_wrapper/vulkan_wrapper.cpp index 53a68fd48..00c9bca34 100644 --- a/test_conformance/common/vulkan_wrapper/vulkan_wrapper.cpp +++ b/test_conformance/common/vulkan_wrapper/vulkan_wrapper.cpp @@ -618,10 +618,15 @@ VulkanQueue::~VulkanQueue() {} void VulkanQueue::submit(const VulkanCommandBuffer &commandBuffer, const VkFence &fence) { + 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, fence); } diff --git a/test_conformance/vulkan/test_vulkan_interop_buffer.cpp b/test_conformance/vulkan/test_vulkan_interop_buffer.cpp index f623f8c0b..e983f209f 100644 --- a/test_conformance/vulkan/test_vulkan_interop_buffer.cpp +++ b/test_conformance/vulkan/test_vulkan_interop_buffer.cpp @@ -275,6 +275,7 @@ int run_test_with_two_queue(cl_context &context, cl_command_queue &cmd_queue1, if (use_fence) { + vkResetFences(vkDevice, 1, &fence); vkWaitForFences(vkDevice, 1, &fence, VK_TRUE, UINT64_MAX); } else @@ -335,6 +336,7 @@ int run_test_with_two_queue(cl_context &context, cl_command_queue &cmd_queue1, { if (use_fence) { + vkResetFences(vkDevice, 1, &fence); vkWaitForFences(vkDevice, 1, &fence, VK_TRUE, UINT64_MAX); } @@ -630,6 +632,7 @@ int run_test_with_one_queue(cl_context &context, cl_command_queue &cmd_queue1, if (use_fence) { + vkResetFences(vkDevice, 1, &fence); vkWaitForFences(vkDevice, 1, &fence, VK_TRUE, UINT64_MAX); } else @@ -666,6 +669,7 @@ int run_test_with_one_queue(cl_context &context, cl_command_queue &cmd_queue1, { if (use_fence) { + vkResetFences(vkDevice, 1, &fence); vkWaitForFences(vkDevice, 1, &fence, VK_TRUE, UINT64_MAX); clFinish(cmd_queue1); @@ -994,6 +998,7 @@ int run_test_with_multi_import_same_ctx( if (use_fence) { + vkResetFences(vkDevice, 1, &fence); vkWaitForFences(vkDevice, 1, &fence, VK_TRUE, UINT64_MAX); } @@ -1038,6 +1043,7 @@ int run_test_with_multi_import_same_ctx( { if (use_fence) { + vkResetFences(vkDevice, 1, &fence); vkWaitForFences(vkDevice, 1, &fence, VK_TRUE, UINT64_MAX); } @@ -1415,6 +1421,7 @@ int run_test_with_multi_import_diff_ctx( if (use_fence) { + vkResetFences(vkDevice, 1, &fence); vkWaitForFences(vkDevice, 1, &fence, VK_TRUE, UINT64_MAX); } @@ -1459,6 +1466,7 @@ int run_test_with_multi_import_diff_ctx( { if (use_fence) { + vkResetFences(vkDevice, 1, &fence); vkWaitForFences(vkDevice, 1, &fence, VK_TRUE, UINT64_MAX); } @@ -1540,6 +1548,7 @@ int run_test_with_multi_import_diff_ctx( { if (use_fence) { + vkResetFences(vkDevice, 1, &fence); vkWaitForFences(vkDevice, 1, &fence, VK_TRUE, UINT64_MAX); } From c5e6dd9330c9654fa31bd5cfddab80b61590d846 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Pawe=C5=82=20Jastrz=C4=99bski?= Date: Wed, 22 Mar 2023 18:25:27 +0100 Subject: [PATCH 07/14] Add Vulkan wrapper for fence. MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Signed-off-by: Paweł Jastrzębski --- .../common/vulkan_wrapper/vulkan_wrapper.cpp | 35 +++++- .../common/vulkan_wrapper/vulkan_wrapper.hpp | 19 ++- .../vulkan/test_vulkan_interop_buffer.cpp | 112 +++++------------- 3 files changed, 77 insertions(+), 89 deletions(-) diff --git a/test_conformance/common/vulkan_wrapper/vulkan_wrapper.cpp b/test_conformance/common/vulkan_wrapper/vulkan_wrapper.cpp index 00c9bca34..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 // //////////////////////////////// @@ -616,7 +647,7 @@ VulkanQueue::VulkanQueue(VkQueue vkQueue): m_vkQueue(vkQueue) {} VulkanQueue::~VulkanQueue() {} void VulkanQueue::submit(const VulkanCommandBuffer &commandBuffer, - const VkFence &fence) + const std::shared_ptr &vkFence) { VulkanCommandBufferList commandBufferList; commandBufferList.add(commandBuffer); @@ -628,7 +659,7 @@ void VulkanQueue::submit(const VulkanCommandBuffer &commandBuffer, vkSubmitInfo.commandBufferCount = (uint32_t)commandBufferList.size(); vkSubmitInfo.pCommandBuffers = commandBufferList(); - vkQueueSubmit(m_vkQueue, 1, &vkSubmitInfo, fence); + vkQueueSubmit(m_vkQueue, 1, &vkSubmitInfo, vkFence->fence); } void VulkanQueue::submit(const VulkanSemaphoreList &waitSemaphoreList, diff --git a/test_conformance/common/vulkan_wrapper/vulkan_wrapper.hpp b/test_conformance/common/vulkan_wrapper/vulkan_wrapper.hpp index b22023e5f..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,7 +172,8 @@ class VulkanQueue { public: const VulkanQueueFamily &getQueueFamily(); - void submit(const VulkanCommandBuffer &commandBuffer, const VkFence &fence); + void submit(const VulkanCommandBuffer &commandBuffer, + const std::shared_ptr &fence); void submit(const VulkanSemaphoreList &waitSemaphoreList, const VulkanCommandBufferList &commandBufferList, const VulkanSemaphoreList &signalSemaphoreList); @@ -570,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/test_vulkan_interop_buffer.cpp b/test_conformance/vulkan/test_vulkan_interop_buffer.cpp index e983f209f..6c302e8f6 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" @@ -118,7 +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); - VkFence fence; + std::shared_ptr fence; VulkanQueue &vkQueue = vkDevice.getQueue(); @@ -140,20 +141,7 @@ int run_test_with_two_queue(cl_context &context, cl_command_queue &cmd_queue1, if (use_fence) { - VkFenceCreateInfo fenceInfo{}; - fenceInfo.sType = VK_STRUCTURE_TYPE_FENCE_CREATE_INFO; - fenceInfo.pNext = nullptr; - fenceInfo.flags = 0; - - VkResult vkStatus = - vkCreateFence(vkDevice, &fenceInfo, nullptr, &fence); - - if (vkStatus != VK_SUCCESS) - { - print_error(vkStatus, "Error: Failed create fence.\n"); - vkDestroyFence(vkDevice, fence, nullptr); - return TEST_FAIL; - } + fence = std::make_shared(vkDevice); } else { @@ -275,8 +263,8 @@ int run_test_with_two_queue(cl_context &context, cl_command_queue &cmd_queue1, if (use_fence) { - vkResetFences(vkDevice, 1, &fence); - vkWaitForFences(vkDevice, 1, &fence, VK_TRUE, UINT64_MAX); + fence->reset(); + fence->wait(); } else { @@ -336,9 +324,8 @@ int run_test_with_two_queue(cl_context &context, cl_command_queue &cmd_queue1, { if (use_fence) { - vkResetFences(vkDevice, 1, &fence); - vkWaitForFences(vkDevice, 1, &fence, VK_TRUE, - UINT64_MAX); + fence->reset(); + fence->wait(); } else { @@ -474,7 +461,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); - VkFence fence; + std::shared_ptr fence; VulkanQueue &vkQueue = vkDevice.getQueue(); @@ -495,20 +482,7 @@ int run_test_with_one_queue(cl_context &context, cl_command_queue &cmd_queue1, if (use_fence) { - VkFenceCreateInfo fenceInfo{}; - fenceInfo.sType = VK_STRUCTURE_TYPE_FENCE_CREATE_INFO; - fenceInfo.pNext = nullptr; - fenceInfo.flags = 0; - - VkResult vkStatus = - vkCreateFence(vkDevice, &fenceInfo, nullptr, &fence); - - if (vkStatus != VK_SUCCESS) - { - print_error(vkStatus, "Error: Failed create fence.\n"); - vkDestroyFence(vkDevice, fence, nullptr); - return TEST_FAIL; - } + fence = std::make_shared(vkDevice); } else { @@ -632,8 +606,8 @@ int run_test_with_one_queue(cl_context &context, cl_command_queue &cmd_queue1, if (use_fence) { - vkResetFences(vkDevice, 1, &fence); - vkWaitForFences(vkDevice, 1, &fence, VK_TRUE, UINT64_MAX); + fence->reset(); + fence->wait(); } else { @@ -669,9 +643,8 @@ int run_test_with_one_queue(cl_context &context, cl_command_queue &cmd_queue1, { if (use_fence) { - vkResetFences(vkDevice, 1, &fence); - vkWaitForFences(vkDevice, 1, &fence, VK_TRUE, - UINT64_MAX); + fence->reset(); + fence->wait(); clFinish(cmd_queue1); } else @@ -805,7 +778,7 @@ int run_test_with_multi_import_same_ctx( getSupportedVulkanExternalSemaphoreHandleTypeList()[0]; VulkanSemaphore vkVk2CLSemaphore(vkDevice, vkExternalSemaphoreHandleType); VulkanSemaphore vkCl2VkSemaphore(vkDevice, vkExternalSemaphoreHandleType); - VkFence fence; + std::shared_ptr fence; VulkanQueue &vkQueue = vkDevice.getQueue(); @@ -827,20 +800,7 @@ int run_test_with_multi_import_same_ctx( if (use_fence) { - VkFenceCreateInfo fenceInfo{}; - fenceInfo.sType = VK_STRUCTURE_TYPE_FENCE_CREATE_INFO; - fenceInfo.pNext = nullptr; - fenceInfo.flags = 0; - - VkResult vkStatus = - vkCreateFence(vkDevice, &fenceInfo, nullptr, &fence); - - if (vkStatus != VK_SUCCESS) - { - print_error(vkStatus, "Error: Failed create fence.\n"); - vkDestroyFence(vkDevice, fence, nullptr); - return TEST_FAIL; - } + fence = std::make_shared(vkDevice); } else { @@ -998,9 +958,8 @@ int run_test_with_multi_import_same_ctx( if (use_fence) { - vkResetFences(vkDevice, 1, &fence); - vkWaitForFences(vkDevice, 1, &fence, VK_TRUE, - UINT64_MAX); + fence->reset(); + fence->wait(); } else { @@ -1043,9 +1002,8 @@ int run_test_with_multi_import_same_ctx( { if (use_fence) { - vkResetFences(vkDevice, 1, &fence); - vkWaitForFences(vkDevice, 1, &fence, VK_TRUE, - UINT64_MAX); + fence->reset(); + fence->wait(); } else { @@ -1203,7 +1161,7 @@ int run_test_with_multi_import_diff_ctx( getSupportedVulkanExternalSemaphoreHandleTypeList()[0]; VulkanSemaphore vkVk2CLSemaphore(vkDevice, vkExternalSemaphoreHandleType); VulkanSemaphore vkCl2VkSemaphore(vkDevice, vkExternalSemaphoreHandleType); - VkFence fence; + std::shared_ptr fence; VulkanQueue &vkQueue = vkDevice.getQueue(); @@ -1225,20 +1183,7 @@ int run_test_with_multi_import_diff_ctx( if (use_fence) { - VkFenceCreateInfo fenceInfo{}; - fenceInfo.sType = VK_STRUCTURE_TYPE_FENCE_CREATE_INFO; - fenceInfo.pNext = nullptr; - fenceInfo.flags = 0; - - VkResult vkStatus = - vkCreateFence(vkDevice, &fenceInfo, nullptr, &fence); - - if (vkStatus != VK_SUCCESS) - { - print_error(vkStatus, "Error: Failed create fence.\n"); - vkDestroyFence(vkDevice, fence, nullptr); - return TEST_FAIL; - } + fence = std::make_shared(vkDevice); } else { @@ -1421,9 +1366,8 @@ int run_test_with_multi_import_diff_ctx( if (use_fence) { - vkResetFences(vkDevice, 1, &fence); - vkWaitForFences(vkDevice, 1, &fence, VK_TRUE, - UINT64_MAX); + fence->reset(); + fence->wait(); } else { @@ -1466,9 +1410,8 @@ int run_test_with_multi_import_diff_ctx( { if (use_fence) { - vkResetFences(vkDevice, 1, &fence); - vkWaitForFences(vkDevice, 1, &fence, VK_TRUE, - UINT64_MAX); + fence->reset(); + fence->wait(); } else { @@ -1548,9 +1491,8 @@ int run_test_with_multi_import_diff_ctx( { if (use_fence) { - vkResetFences(vkDevice, 1, &fence); - vkWaitForFences(vkDevice, 1, &fence, VK_TRUE, - UINT64_MAX); + fence->reset(); + fence->wait(); } else { From 8612f4a57757890da0506ca399cb5542966c67cc Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Pawe=C5=82=20Jastrz=C4=99bski?= Date: Mon, 17 Apr 2023 15:46:49 +0200 Subject: [PATCH 08/14] Rework fence reset. MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Signed-off-by: Paweł Jastrzębski --- .../vulkan/test_vulkan_interop_buffer.cpp | 19 ++++++++++--------- 1 file changed, 10 insertions(+), 9 deletions(-) diff --git a/test_conformance/vulkan/test_vulkan_interop_buffer.cpp b/test_conformance/vulkan/test_vulkan_interop_buffer.cpp index 6c302e8f6..4804293ba 100644 --- a/test_conformance/vulkan/test_vulkan_interop_buffer.cpp +++ b/test_conformance/vulkan/test_vulkan_interop_buffer.cpp @@ -241,6 +241,7 @@ int run_test_with_two_queue(cl_context &context, cl_command_queue &cmd_queue1, { if (use_fence) { + fence->reset(); vkQueue.submit(vkCommandBuffer, fence); } else @@ -252,6 +253,7 @@ int run_test_with_two_queue(cl_context &context, cl_command_queue &cmd_queue1, { if (use_fence) { + fence->reset(); vkQueue.submit(vkCommandBuffer, fence); } else @@ -263,7 +265,6 @@ int run_test_with_two_queue(cl_context &context, cl_command_queue &cmd_queue1, if (use_fence) { - fence->reset(); fence->wait(); } else @@ -324,7 +325,6 @@ int run_test_with_two_queue(cl_context &context, cl_command_queue &cmd_queue1, { if (use_fence) { - fence->reset(); fence->wait(); } else @@ -584,6 +584,7 @@ int run_test_with_one_queue(cl_context &context, cl_command_queue &cmd_queue1, { if (use_fence) { + fence->reset(); vkQueue.submit(vkCommandBuffer, fence); } else @@ -595,6 +596,7 @@ int run_test_with_one_queue(cl_context &context, cl_command_queue &cmd_queue1, { if (use_fence) { + fence->reset(); vkQueue.submit(vkCommandBuffer, fence); } else @@ -606,7 +608,6 @@ int run_test_with_one_queue(cl_context &context, cl_command_queue &cmd_queue1, if (use_fence) { - fence->reset(); fence->wait(); } else @@ -643,7 +644,6 @@ int run_test_with_one_queue(cl_context &context, cl_command_queue &cmd_queue1, { if (use_fence) { - fence->reset(); fence->wait(); clFinish(cmd_queue1); } @@ -936,6 +936,7 @@ int run_test_with_multi_import_same_ctx( { if (use_fence) { + fence->reset(); vkQueue.submit(vkCommandBuffer, fence); } else @@ -947,6 +948,7 @@ int run_test_with_multi_import_same_ctx( { if (use_fence) { + fence->reset(); vkQueue.submit(vkCommandBuffer, fence); } else @@ -958,7 +960,6 @@ int run_test_with_multi_import_same_ctx( if (use_fence) { - fence->reset(); fence->wait(); } else @@ -1002,7 +1003,6 @@ int run_test_with_multi_import_same_ctx( { if (use_fence) { - fence->reset(); fence->wait(); } else @@ -1344,6 +1344,7 @@ int run_test_with_multi_import_diff_ctx( { if (use_fence) { + fence->reset(); vkQueue.submit(vkCommandBuffer, fence); } else @@ -1355,6 +1356,7 @@ int run_test_with_multi_import_diff_ctx( { if (use_fence) { + fence->reset(); vkQueue.submit(vkCommandBuffer, fence); } else @@ -1366,7 +1368,6 @@ int run_test_with_multi_import_diff_ctx( if (use_fence) { - fence->reset(); fence->wait(); } else @@ -1410,7 +1411,6 @@ int run_test_with_multi_import_diff_ctx( { if (use_fence) { - fence->reset(); fence->wait(); } else @@ -1426,6 +1426,7 @@ int run_test_with_multi_import_diff_ctx( { if (use_fence) { + fence->reset(); vkQueue.submit(vkCommandBuffer, fence); } else @@ -1437,6 +1438,7 @@ int run_test_with_multi_import_diff_ctx( { if (use_fence) { + fence->reset(); vkQueue.submit(vkCommandBuffer, fence); } else @@ -1491,7 +1493,6 @@ int run_test_with_multi_import_diff_ctx( { if (use_fence) { - fence->reset(); fence->wait(); } else From 51cc730532bb3069ff3b5a24ebfb162c6869da5d Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Pawe=C5=82=20Jastrz=C4=99bski?= Date: Wed, 17 May 2023 15:36:55 +0200 Subject: [PATCH 09/14] Change synchronisation mechanisms. MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Changes made: - wait for fence with clFinish - queue submit with wait for fence Signed-off-by: Paweł Jastrzębski --- .../vulkan/test_vulkan_interop_buffer.cpp | 19 +++++++++---------- 1 file changed, 9 insertions(+), 10 deletions(-) diff --git a/test_conformance/vulkan/test_vulkan_interop_buffer.cpp b/test_conformance/vulkan/test_vulkan_interop_buffer.cpp index 4804293ba..3ae9a2393 100644 --- a/test_conformance/vulkan/test_vulkan_interop_buffer.cpp +++ b/test_conformance/vulkan/test_vulkan_interop_buffer.cpp @@ -265,7 +265,7 @@ int run_test_with_two_queue(cl_context &context, cl_command_queue &cmd_queue1, if (use_fence) { - fence->wait(); + clFinish(cmd_queue1); } else { @@ -325,7 +325,7 @@ int run_test_with_two_queue(cl_context &context, cl_command_queue &cmd_queue1, { if (use_fence) { - fence->wait(); + clFinish(cmd_queue2); } else { @@ -608,7 +608,7 @@ int run_test_with_one_queue(cl_context &context, cl_command_queue &cmd_queue1, if (use_fence) { - fence->wait(); + clFinish(cmd_queue1); } else { @@ -644,7 +644,6 @@ int run_test_with_one_queue(cl_context &context, cl_command_queue &cmd_queue1, { if (use_fence) { - fence->wait(); clFinish(cmd_queue1); } else @@ -960,7 +959,7 @@ int run_test_with_multi_import_same_ctx( if (use_fence) { - fence->wait(); + clFinish(cmd_queue1); } else { @@ -1003,7 +1002,7 @@ int run_test_with_multi_import_same_ctx( { if (use_fence) { - fence->wait(); + clFinish(cmd_queue1); } else { @@ -1368,7 +1367,7 @@ int run_test_with_multi_import_diff_ctx( if (use_fence) { - fence->wait(); + clFinish(cmd_queue1); } else { @@ -1411,7 +1410,7 @@ int run_test_with_multi_import_diff_ctx( { if (use_fence) { - fence->wait(); + clFinish(cmd_queue1); } else { @@ -1450,7 +1449,7 @@ int run_test_with_multi_import_diff_ctx( if (use_fence) { - vkQueue.submit(vkCommandBuffer, fence); + fence->wait(); } else { @@ -1493,7 +1492,7 @@ int run_test_with_multi_import_diff_ctx( { if (use_fence) { - fence->wait(); + clFinish(cmd_queue2); } else { From eb69cecfb56b78e362428cf155f14982560ff619 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Pawe=C5=82=20Jastrz=C4=99bski?= Date: Mon, 22 May 2023 14:31:06 +0200 Subject: [PATCH 10/14] Replace clFinish with vkWaitForFences. MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Replaced clFinish with vkWaitForFences in Vulkan exectution context. Signed-off-by: Paweł Jastrzębski --- test_conformance/vulkan/test_vulkan_interop_buffer.cpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/test_conformance/vulkan/test_vulkan_interop_buffer.cpp b/test_conformance/vulkan/test_vulkan_interop_buffer.cpp index 3ae9a2393..982c9c23d 100644 --- a/test_conformance/vulkan/test_vulkan_interop_buffer.cpp +++ b/test_conformance/vulkan/test_vulkan_interop_buffer.cpp @@ -265,7 +265,7 @@ int run_test_with_two_queue(cl_context &context, cl_command_queue &cmd_queue1, if (use_fence) { - clFinish(cmd_queue1); + fence->wait(); } else { @@ -608,7 +608,7 @@ int run_test_with_one_queue(cl_context &context, cl_command_queue &cmd_queue1, if (use_fence) { - clFinish(cmd_queue1); + fence->wait(); } else { From 24dd4897c12a5c4a5a330bb89d5dbe4beeb56085 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Pawe=C5=82=20Jastrz=C4=99bski?= Date: Tue, 23 May 2023 11:32:17 +0200 Subject: [PATCH 11/14] Replace remaining clFinish with vkWaitForFences. MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Replaced remaining clFinish with vkWaitForFences in Vulkan exectution context. Signed-off-by: Paweł Jastrzębski --- test_conformance/vulkan/test_vulkan_interop_buffer.cpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/test_conformance/vulkan/test_vulkan_interop_buffer.cpp b/test_conformance/vulkan/test_vulkan_interop_buffer.cpp index 982c9c23d..070959287 100644 --- a/test_conformance/vulkan/test_vulkan_interop_buffer.cpp +++ b/test_conformance/vulkan/test_vulkan_interop_buffer.cpp @@ -959,7 +959,7 @@ int run_test_with_multi_import_same_ctx( if (use_fence) { - clFinish(cmd_queue1); + fence->wait(); } else { @@ -1367,7 +1367,7 @@ int run_test_with_multi_import_diff_ctx( if (use_fence) { - clFinish(cmd_queue1); + fence->wait(); } else { From a5f1df7abed623ffe2b2fab8e00d6ce9d7ddd339 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Pawe=C5=82=20Jastrz=C4=99bski?= Date: Wed, 21 Jun 2023 16:08:27 +0200 Subject: [PATCH 12/14] Fix review comments for synchoronisation simplification. MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Signed-off-by: Paweł Jastrzębski --- .../vulkan/test_vulkan_interop_buffer.cpp | 94 +++++++------------ 1 file changed, 32 insertions(+), 62 deletions(-) diff --git a/test_conformance/vulkan/test_vulkan_interop_buffer.cpp b/test_conformance/vulkan/test_vulkan_interop_buffer.cpp index 070959287..b1eff5002 100644 --- a/test_conformance/vulkan/test_vulkan_interop_buffer.cpp +++ b/test_conformance/vulkan/test_vulkan_interop_buffer.cpp @@ -119,7 +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; + std::shared_ptr fence = nullptr; VulkanQueue &vkQueue = vkDevice.getQueue(); @@ -237,38 +237,24 @@ 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) { - if (use_fence) - { - fence->reset(); - vkQueue.submit(vkCommandBuffer, fence); - } - else - { - vkQueue.submit(vkCommandBuffer, vkVk2CLSemaphore); - } + fence->reset(); + vkQueue.submit(vkCommandBuffer, fence); + fence->wait(); } else { - if (use_fence) + if (iter == 0) { - fence->reset(); - vkQueue.submit(vkCommandBuffer, fence); + vkQueue.submit(vkCommandBuffer, vkVk2CLSemaphore); } else { vkQueue.submit(vkCl2VkSemaphore, vkCommandBuffer, vkVk2CLSemaphore); } - } - if (use_fence) - { - fence->wait(); - } - else - { clVk2CLExternalSemaphore->wait(cmd_queue1); } @@ -321,16 +307,16 @@ int run_test_with_two_queue(cl_context &context, cl_command_queue &cmd_queue1, goto CLEANUP; } - if (iter != (maxIter - 1)) + if (use_fence) { - if (use_fence) - { - clFinish(cmd_queue2); - } - else - { - clCl2VkExternalSemaphore->signal(cmd_queue2); - } + clFlush(cmd_queue1); + clFlush(cmd_queue2); + clFinish(cmd_queue1); + clFinish(cmd_queue2); + } + else if (!use_fence && iter != (maxIter - 1)) + { + clCl2VkExternalSemaphore->signal(cmd_queue2); } } error_2 = (uint8_t *)malloc(sizeof(uint8_t)); @@ -461,7 +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; + std::shared_ptr fence = nullptr; VulkanQueue &vkQueue = vkDevice.getQueue(); @@ -580,38 +566,24 @@ 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) { - if (use_fence) - { - fence->reset(); - vkQueue.submit(vkCommandBuffer, fence); - } - else - { - vkQueue.submit(vkCommandBuffer, vkVk2CLSemaphore); - } + fence->reset(); + vkQueue.submit(vkCommandBuffer, fence); + fence->wait(); } else { - if (use_fence) + if (iter == 0) { - fence->reset(); - vkQueue.submit(vkCommandBuffer, fence); + vkQueue.submit(vkCommandBuffer, vkVk2CLSemaphore); } else { vkQueue.submit(vkCl2VkSemaphore, vkCommandBuffer, vkVk2CLSemaphore); } - } - if (use_fence) - { - fence->wait(); - } - else - { clVk2CLExternalSemaphore->wait(cmd_queue1); } @@ -640,16 +612,14 @@ 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) { - if (use_fence) - { - clFinish(cmd_queue1); - } - else - { - clCl2VkExternalSemaphore->signal(cmd_queue1); - } + clFlush(cmd_queue1); + clFinish(cmd_queue1); + } + else if (!use_fence && (iter != (maxIter - 1))) + { + clCl2VkExternalSemaphore->signal(cmd_queue1); } } error_2 = (uint8_t *)malloc(sizeof(uint8_t)); @@ -777,7 +747,7 @@ int run_test_with_multi_import_same_ctx( getSupportedVulkanExternalSemaphoreHandleTypeList()[0]; VulkanSemaphore vkVk2CLSemaphore(vkDevice, vkExternalSemaphoreHandleType); VulkanSemaphore vkCl2VkSemaphore(vkDevice, vkExternalSemaphoreHandleType); - std::shared_ptr fence; + std::shared_ptr fence = nullptr; VulkanQueue &vkQueue = vkDevice.getQueue(); @@ -1160,7 +1130,7 @@ int run_test_with_multi_import_diff_ctx( getSupportedVulkanExternalSemaphoreHandleTypeList()[0]; VulkanSemaphore vkVk2CLSemaphore(vkDevice, vkExternalSemaphoreHandleType); VulkanSemaphore vkCl2VkSemaphore(vkDevice, vkExternalSemaphoreHandleType); - std::shared_ptr fence; + std::shared_ptr fence = nullptr; VulkanQueue &vkQueue = vkDevice.getQueue(); From 6436ec7a1e1c202e5fefe5b55721c27ce780c243 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Pawe=C5=82=20Jastrz=C4=99bski?= Date: Fri, 23 Jun 2023 13:29:13 +0200 Subject: [PATCH 13/14] Fix review comments for synchoronisation simplification for remaining tests. MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Signed-off-by: Paweł Jastrzębski --- .../vulkan/test_vulkan_interop_buffer.cpp | 102 +++++++----------- 1 file changed, 36 insertions(+), 66 deletions(-) diff --git a/test_conformance/vulkan/test_vulkan_interop_buffer.cpp b/test_conformance/vulkan/test_vulkan_interop_buffer.cpp index b1eff5002..c141bbaee 100644 --- a/test_conformance/vulkan/test_vulkan_interop_buffer.cpp +++ b/test_conformance/vulkan/test_vulkan_interop_buffer.cpp @@ -901,24 +901,17 @@ int run_test_with_multi_import_same_ctx( for (uint32_t iter = 0; iter < maxIter; iter++) { - if (iter == 0) + if (use_fence) { - if (use_fence) - { - fence->reset(); - vkQueue.submit(vkCommandBuffer, fence); - } - else - { - vkQueue.submit(vkCommandBuffer, vkVk2CLSemaphore); - } + fence->reset(); + vkQueue.submit(vkCommandBuffer, fence); + fence->wait(); } else { - if (use_fence) + if (iter == 0) { - fence->reset(); - vkQueue.submit(vkCommandBuffer, fence); + vkQueue.submit(vkCommandBuffer, vkVk2CLSemaphore); } else { @@ -968,16 +961,13 @@ int run_test_with_multi_import_same_ctx( goto CLEANUP; } } - if (iter != (maxIter - 1)) + if (use_fence) { - if (use_fence) - { - clFinish(cmd_queue1); - } - else - { - clCl2VkExternalSemaphore->signal(cmd_queue1); - } + clFinish(cmd_queue1); + } + else if (!!use_fence && iter != (maxIter - 1)) + { + clCl2VkExternalSemaphore->signal(cmd_queue1); } } error_2 = (uint8_t *)malloc(sizeof(uint8_t)); @@ -1309,24 +1299,17 @@ int run_test_with_multi_import_diff_ctx( for (uint32_t iter = 0; iter < maxIter; iter++) { - if (iter == 0) + if (use_fence) { - if (use_fence) - { - fence->reset(); - vkQueue.submit(vkCommandBuffer, fence); - } - else - { - vkQueue.submit(vkCommandBuffer, vkVk2CLSemaphore); - } + fence->reset(); + vkQueue.submit(vkCommandBuffer, fence); + fence->wait(); } else { - if (use_fence) + if (iter == 0) { - fence->reset(); - vkQueue.submit(vkCommandBuffer, fence); + vkQueue.submit(vkCommandBuffer, vkVk2CLSemaphore); } else { @@ -1376,39 +1359,29 @@ int run_test_with_multi_import_diff_ctx( goto CLEANUP; } } - if (iter != (maxIter - 1)) + if (use_fence) { - if (use_fence) - { - clFinish(cmd_queue1); - } - else - { - clCl2VkExternalSemaphore->signal(cmd_queue1); - } + clFinish(cmd_queue1); + } + else if (!use_fence && iter != (maxIter - 1)) + { + clCl2VkExternalSemaphore->signal(cmd_queue1); } } clFinish(cmd_queue1); for (uint32_t iter = 0; iter < maxIter; iter++) { - if (iter == 0) + if (use_fence) { - if (use_fence) - { - fence->reset(); - vkQueue.submit(vkCommandBuffer, fence); - } - else - { - vkQueue.submit(vkCommandBuffer, vkVk2CLSemaphore); - } + fence->reset(); + vkQueue.submit(vkCommandBuffer, fence); + fence->wait(); } else { - if (use_fence) + if (iter == 0) { - fence->reset(); - vkQueue.submit(vkCommandBuffer, fence); + vkQueue.submit(vkCommandBuffer, vkVk2CLSemaphore); } else { @@ -1458,16 +1431,13 @@ int run_test_with_multi_import_diff_ctx( goto CLEANUP; } } - if (iter != (maxIter - 1)) + if (use_fence) { - if (use_fence) - { - clFinish(cmd_queue2); - } - else - { - clCl2VkExternalSemaphore2->signal(cmd_queue2); - } + clFinish(cmd_queue2); + } + else if (!use_fence && iter != (maxIter - 1)) + { + clCl2VkExternalSemaphore2->signal(cmd_queue2); } } clFinish(cmd_queue2); From 3e9f0a7e258a5ea11c1c7a4b9bb4013d8200eaeb Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Pawe=C5=82=20Jastrz=C4=99bski?= Date: Tue, 27 Jun 2023 13:15:30 +0200 Subject: [PATCH 14/14] Fix condition check. MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Signed-off-by: Paweł Jastrzębski --- test_conformance/vulkan/test_vulkan_interop_buffer.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/test_conformance/vulkan/test_vulkan_interop_buffer.cpp b/test_conformance/vulkan/test_vulkan_interop_buffer.cpp index c141bbaee..5390ef690 100644 --- a/test_conformance/vulkan/test_vulkan_interop_buffer.cpp +++ b/test_conformance/vulkan/test_vulkan_interop_buffer.cpp @@ -965,7 +965,7 @@ int run_test_with_multi_import_same_ctx( { clFinish(cmd_queue1); } - else if (!!use_fence && iter != (maxIter - 1)) + else if (!use_fence && iter != (maxIter - 1)) { clCl2VkExternalSemaphore->signal(cmd_queue1); }