diff --git a/test_conformance/vulkan/main.cpp b/test_conformance/vulkan/main.cpp index 6cbde5cc..2eeb0c36 100644 --- a/test_conformance/vulkan/main.cpp +++ b/test_conformance/vulkan/main.cpp @@ -134,7 +134,6 @@ cl_device_id *devices; const size_t bufsize = BUFFERSIZE; char buf[BUFFERSIZE]; cl_uchar uuid[CL_UUID_SIZE_KHR]; -VulkanDevice vkDevice; unsigned int numCQ; bool multiImport; bool multiCtx; @@ -220,9 +219,12 @@ int main(int argc, const char *argv[]) if (!checkVkSupport()) { log_info("Vulkan supported GPU not found \n"); + log_info("TEST SKIPPED \n"); return 0; } + VulkanDevice vkDevice; + cl_device_type requestedDeviceType = CL_DEVICE_TYPE_GPU; char *force_cpu = getenv("CL_DEVICE_TYPE"); if (force_cpu != NULL) diff --git a/test_conformance/vulkan/shaders/buffer.comp b/test_conformance/vulkan/shaders/buffer.comp new file mode 100644 index 00000000..d8756f92 --- /dev/null +++ b/test_conformance/vulkan/shaders/buffer.comp @@ -0,0 +1,28 @@ +#version 450 +#extension GL_ARB_separate_shader_objects : enable +#extension GL_EXT_shader_explicit_arithmetic_types_int8 : enable +#extension GL_EXT_shader_explicit_arithmetic_types_int32 : enable + +#define MAX_BUFFERS 5 + +layout(binding = 0) buffer Params +{ + uint32_t numBuffers; + uint32_t bufferSize; + uint32_t interBufferOffset; +}; +layout(binding = 1) buffer Buffer +{ + uint8_t ptr[]; +} bufferPtrList[MAX_BUFFERS]; +layout(local_size_x = 512) in; +void main() { + for (uint32_t bufIdx = 0; bufIdx < numBuffers; bufIdx++) { + uint32_t ptrIdx = gl_GlobalInvocationID.x; + uint32_t limit = bufferSize; + while (ptrIdx < limit) { + bufferPtrList[bufIdx].ptr[ptrIdx]++; + ptrIdx += (gl_NumWorkGroups.x * gl_WorkGroupSize.x); + } + } +} \ No newline at end of file diff --git a/test_conformance/vulkan/shaders/buffer.spv b/test_conformance/vulkan/shaders/buffer.spv new file mode 100644 index 00000000..685523ba Binary files /dev/null and b/test_conformance/vulkan/shaders/buffer.spv differ diff --git a/test_conformance/vulkan/shaders/image2D.comp b/test_conformance/vulkan/shaders/image2D.comp new file mode 100644 index 00000000..42fa2f73 --- /dev/null +++ b/test_conformance/vulkan/shaders/image2D.comp @@ -0,0 +1,31 @@ +#version 450 +#extension GL_ARB_separate_shader_objects : enable +#extension GL_EXT_shader_explicit_arithmetic_types_int32 : enable + +#define MAX_2D_IMAGES 5 +#define MAX_2D_IMAGE_MIP_LEVELS 11 +#define MAX_2D_IMAGE_DESCRIPTORS MAX_2D_IMAGES * MAX_2D_IMAGE_MIP_LEVELS + +layout(binding = 0) buffer Params +{ + uint32_t numImage2DDescriptors; +}; +layout(binding = 1, rgba32f ) uniform image2D image2DList[ MAX_2D_IMAGE_DESCRIPTORS ]; +layout(local_size_x = 32, local_size_y = 32) in; +void main() { + uvec3 numThreads = gl_NumWorkGroups * gl_WorkGroupSize; + for (uint32_t image2DIdx = 0; image2DIdx < numImage2DDescriptors; image2DIdx++) { + ivec2 imageDim = imageSize(image2DList[image2DIdx]); + uint32_t heightBy2 = imageDim.y / 2; + for (uint32_t row = gl_GlobalInvocationID.y; row < heightBy2; row += numThreads.y) { + for (uint32_t col = gl_GlobalInvocationID.x; col < imageDim.x; col += numThreads.x) { + ivec2 coordsA = ivec2(col, row); + ivec2 coordsB = ivec2(col, imageDim.y - row - 1); + vec4 dataA = imageLoad(image2DList[image2DIdx], coordsA); + vec4 dataB = imageLoad(image2DList[image2DIdx], coordsB); + imageStore(image2DList[image2DIdx], coordsA, dataB); + imageStore(image2DList[image2DIdx], coordsB, dataA); + } + } + } +} \ No newline at end of file diff --git a/test_conformance/vulkan/shaders/image2D_r16i.spv b/test_conformance/vulkan/shaders/image2D_r16i.spv new file mode 100644 index 00000000..00c5c283 Binary files /dev/null and b/test_conformance/vulkan/shaders/image2D_r16i.spv differ diff --git a/test_conformance/vulkan/shaders/image2D_r16ui.spv b/test_conformance/vulkan/shaders/image2D_r16ui.spv new file mode 100644 index 00000000..87514d9f Binary files /dev/null and b/test_conformance/vulkan/shaders/image2D_r16ui.spv differ diff --git a/test_conformance/vulkan/shaders/image2D_r32f.spv b/test_conformance/vulkan/shaders/image2D_r32f.spv new file mode 100644 index 00000000..e82c9c19 Binary files /dev/null and b/test_conformance/vulkan/shaders/image2D_r32f.spv differ diff --git a/test_conformance/vulkan/shaders/image2D_r32i.spv b/test_conformance/vulkan/shaders/image2D_r32i.spv new file mode 100644 index 00000000..7ea8d26f Binary files /dev/null and b/test_conformance/vulkan/shaders/image2D_r32i.spv differ diff --git a/test_conformance/vulkan/shaders/image2D_r32ui.spv b/test_conformance/vulkan/shaders/image2D_r32ui.spv new file mode 100644 index 00000000..dbcdbc5f Binary files /dev/null and b/test_conformance/vulkan/shaders/image2D_r32ui.spv differ diff --git a/test_conformance/vulkan/shaders/image2D_r8i.spv b/test_conformance/vulkan/shaders/image2D_r8i.spv new file mode 100644 index 00000000..1a641475 Binary files /dev/null and b/test_conformance/vulkan/shaders/image2D_r8i.spv differ diff --git a/test_conformance/vulkan/shaders/image2D_r8ui.spv b/test_conformance/vulkan/shaders/image2D_r8ui.spv new file mode 100644 index 00000000..a90ccf98 Binary files /dev/null and b/test_conformance/vulkan/shaders/image2D_r8ui.spv differ diff --git a/test_conformance/vulkan/shaders/image2D_rg16i.spv b/test_conformance/vulkan/shaders/image2D_rg16i.spv new file mode 100644 index 00000000..07996173 Binary files /dev/null and b/test_conformance/vulkan/shaders/image2D_rg16i.spv differ diff --git a/test_conformance/vulkan/shaders/image2D_rg16ui.spv b/test_conformance/vulkan/shaders/image2D_rg16ui.spv new file mode 100644 index 00000000..f73e096b Binary files /dev/null and b/test_conformance/vulkan/shaders/image2D_rg16ui.spv differ diff --git a/test_conformance/vulkan/shaders/image2D_rg32f.spv b/test_conformance/vulkan/shaders/image2D_rg32f.spv new file mode 100644 index 00000000..1489660e Binary files /dev/null and b/test_conformance/vulkan/shaders/image2D_rg32f.spv differ diff --git a/test_conformance/vulkan/shaders/image2D_rg32i.spv b/test_conformance/vulkan/shaders/image2D_rg32i.spv new file mode 100644 index 00000000..b7d302f4 Binary files /dev/null and b/test_conformance/vulkan/shaders/image2D_rg32i.spv differ diff --git a/test_conformance/vulkan/shaders/image2D_rg32ui.spv b/test_conformance/vulkan/shaders/image2D_rg32ui.spv new file mode 100644 index 00000000..6cf2f1b8 Binary files /dev/null and b/test_conformance/vulkan/shaders/image2D_rg32ui.spv differ diff --git a/test_conformance/vulkan/shaders/image2D_rg8i.spv b/test_conformance/vulkan/shaders/image2D_rg8i.spv new file mode 100644 index 00000000..a71b9bf0 Binary files /dev/null and b/test_conformance/vulkan/shaders/image2D_rg8i.spv differ diff --git a/test_conformance/vulkan/shaders/image2D_rg8ui.spv b/test_conformance/vulkan/shaders/image2D_rg8ui.spv new file mode 100644 index 00000000..2aca9290 Binary files /dev/null and b/test_conformance/vulkan/shaders/image2D_rg8ui.spv differ diff --git a/test_conformance/vulkan/shaders/image2D_rgba16i.spv b/test_conformance/vulkan/shaders/image2D_rgba16i.spv new file mode 100644 index 00000000..0cb95dfd Binary files /dev/null and b/test_conformance/vulkan/shaders/image2D_rgba16i.spv differ diff --git a/test_conformance/vulkan/shaders/image2D_rgba16ui.spv b/test_conformance/vulkan/shaders/image2D_rgba16ui.spv new file mode 100644 index 00000000..84c3d3db Binary files /dev/null and b/test_conformance/vulkan/shaders/image2D_rgba16ui.spv differ diff --git a/test_conformance/vulkan/shaders/image2D_rgba32f.spv b/test_conformance/vulkan/shaders/image2D_rgba32f.spv new file mode 100644 index 00000000..35136c58 Binary files /dev/null and b/test_conformance/vulkan/shaders/image2D_rgba32f.spv differ diff --git a/test_conformance/vulkan/shaders/image2D_rgba32i.spv b/test_conformance/vulkan/shaders/image2D_rgba32i.spv new file mode 100644 index 00000000..4d1ae581 Binary files /dev/null and b/test_conformance/vulkan/shaders/image2D_rgba32i.spv differ diff --git a/test_conformance/vulkan/shaders/image2D_rgba32ui.spv b/test_conformance/vulkan/shaders/image2D_rgba32ui.spv new file mode 100644 index 00000000..bed86f0c Binary files /dev/null and b/test_conformance/vulkan/shaders/image2D_rgba32ui.spv differ diff --git a/test_conformance/vulkan/shaders/image2D_rgba8i.spv b/test_conformance/vulkan/shaders/image2D_rgba8i.spv new file mode 100644 index 00000000..edf8c58c Binary files /dev/null and b/test_conformance/vulkan/shaders/image2D_rgba8i.spv differ diff --git a/test_conformance/vulkan/shaders/image2D_rgba8ui.spv b/test_conformance/vulkan/shaders/image2D_rgba8ui.spv new file mode 100644 index 00000000..bb9a770c Binary files /dev/null and b/test_conformance/vulkan/shaders/image2D_rgba8ui.spv differ diff --git a/test_conformance/vulkan/test_vulkan_api_consistency.cpp b/test_conformance/vulkan/test_vulkan_api_consistency.cpp index 2987418f..f22ac319 100644 --- a/test_conformance/vulkan/test_vulkan_api_consistency.cpp +++ b/test_conformance/vulkan/test_vulkan_api_consistency.cpp @@ -238,7 +238,7 @@ int test_consistency_external_image(cl_device_id deviceID, cl_context _context, const VulkanMemoryTypeList& memoryTypeList = vkImage2D->getMemoryTypeList(); uint64_t totalImageMemSize = vkImage2D->getSize(); - log_info("Memory type index: %d\n", (uint32_t)memoryTypeList[0]); + log_info("Memory type index: %lu\n", (uint32_t)memoryTypeList[0]); log_info("Memory type property: %d\n", memoryTypeList[0].getMemoryTypeProperty()); log_info("Image size : %d\n", totalImageMemSize); @@ -552,17 +552,17 @@ int test_consistency_external_semaphore(cl_device_id deviceID, // Pass invalid object to release call - errNum = clReleaseSemaphoreObjectKHRptr(NULL); + errNum = clReleaseSemaphoreKHRptr(NULL); test_failure_error(errNum, CL_INVALID_VALUE, - "clReleaseSemaphoreObjectKHRptr fails with " + "clReleaseSemaphoreKHRptr fails with " "CL_INVALID_VALUE when NULL semaphore object is passed"); // Release both semaphore objects - errNum = clReleaseSemaphoreObjectKHRptr(clVk2Clsemaphore); - test_error(errNum, "clReleaseSemaphoreObjectKHRptr failed"); + errNum = clReleaseSemaphoreKHRptr(clVk2Clsemaphore); + test_error(errNum, "clReleaseSemaphoreKHRptr failed"); - errNum = clReleaseSemaphoreObjectKHRptr(clCl2Vksemaphore); - test_error(errNum, "clReleaseSemaphoreObjectKHRptr failed"); + errNum = clReleaseSemaphoreKHRptr(clCl2Vksemaphore); + test_error(errNum, "clReleaseSemaphoreKHRptr failed"); return TEST_PASS; } diff --git a/test_conformance/vulkan/test_vulkan_interop_buffer.cpp b/test_conformance/vulkan/test_vulkan_interop_buffer.cpp index 7daf96de..9b0bc9de 100644 --- a/test_conformance/vulkan/test_vulkan_interop_buffer.cpp +++ b/test_conformance/vulkan/test_vulkan_interop_buffer.cpp @@ -39,35 +39,6 @@ struct Params }; } -static const char *vkBufferShader = - "#version 450\n" - "#extension GL_ARB_separate_shader_objects : enable\n" - "#extension GL_NV_gpu_shader5 : enable\n" - "layout(binding = 0) buffer Params\n" - "{\n" - " uint32_t numBuffers;\n" - " uint32_t bufferSize;\n" - " uint32_t interBufferOffset;\n" - "};\n" - "layout(binding = 1) buffer Buffer\n" - "{\n" - " uint8_t ptr[];\n" - "} bufferPtrList[" STRING( - MAX_BUFFERS) "];\n" - "layout(local_size_x = 512) in;\n" - "void main() {\n" - " for (uint32_t bufIdx = 0; bufIdx < numBuffers;" - " bufIdx++) {\n" - " uint32_t ptrIdx = gl_GlobalInvocationID.x;\n" - " uint32_t limit = bufferSize;\n" - " while (ptrIdx < limit) {\n" - " bufferPtrList[bufIdx].ptr[ptrIdx]++;\n" - " ptrIdx += (gl_NumWorkGroups.x * " - "gl_WorkGroupSize.x);\n" - " }\n" - " }\n" - "}\n"; - const char *kernel_text_numbuffer_1 = " \ __kernel void clUpdateBuffer(int bufferSize, __global unsigned char *a) { \n\ int gid = get_global_id(0); \n\ @@ -149,6 +120,8 @@ int run_test_with_two_queue(cl_context &context, cl_command_queue &cmd_queue1, VulkanQueue &vkQueue = vkDevice.getQueue(); + std::vector vkBufferShader = readFile("buffer.spv"); + VulkanShaderModule vkBufferShaderModule(vkDevice, vkBufferShader); VulkanDescriptorSetLayoutBindingList vkDescriptorSetLayoutBindingList( MAX_BUFFERS + 1, VULKAN_DESCRIPTOR_TYPE_STORAGE_BUFFER); @@ -446,6 +419,7 @@ int run_test_with_one_queue(cl_context &context, cl_command_queue &cmd_queue1, VulkanQueue &vkQueue = vkDevice.getQueue(); + std::vector vkBufferShader = readFile("buffer.spv"); VulkanShaderModule vkBufferShaderModule(vkDevice, vkBufferShader); VulkanDescriptorSetLayoutBindingList vkDescriptorSetLayoutBindingList( MAX_BUFFERS + 1, VULKAN_DESCRIPTOR_TYPE_STORAGE_BUFFER); @@ -716,6 +690,8 @@ int run_test_with_multi_import_same_ctx( VulkanQueue &vkQueue = vkDevice.getQueue(); + std::vector vkBufferShader = readFile("buffer.spv"); + VulkanShaderModule vkBufferShaderModule(vkDevice, vkBufferShader); VulkanDescriptorSetLayoutBindingList vkDescriptorSetLayoutBindingList( MAX_BUFFERS + 1, VULKAN_DESCRIPTOR_TYPE_STORAGE_BUFFER); @@ -1050,6 +1026,8 @@ int run_test_with_multi_import_diff_ctx( VulkanQueue &vkQueue = vkDevice.getQueue(); + std::vector vkBufferShader = readFile("buffer.spv"); + VulkanShaderModule vkBufferShaderModule(vkDevice, vkBufferShader); VulkanDescriptorSetLayoutBindingList vkDescriptorSetLayoutBindingList( MAX_BUFFERS + 1, VULKAN_DESCRIPTOR_TYPE_STORAGE_BUFFER); diff --git a/test_conformance/vulkan/test_vulkan_interop_image.cpp b/test_conformance/vulkan/test_vulkan_interop_image.cpp index f1d0af1f..7577de09 100644 --- a/test_conformance/vulkan/test_vulkan_interop_image.cpp +++ b/test_conformance/vulkan/test_vulkan_interop_image.cpp @@ -25,8 +25,6 @@ #define MAX_2D_IMAGE_ELEMENT_SIZE 16 #define MAX_2D_IMAGE_MIP_LEVELS 11 #define MAX_2D_IMAGE_DESCRIPTORS MAX_2D_IMAGES *MAX_2D_IMAGE_MIP_LEVELS -#define GLSL_FORMAT_STRING "" -#define GLSL_TYPE_PREFIX_STRING "" #define NUM_THREADS_PER_GROUP_X 32 #define NUM_THREADS_PER_GROUP_Y 32 #define NUM_BLOCKS(size, blockSize) \ @@ -54,61 +52,8 @@ struct Params } static cl_uchar uuid[CL_UUID_SIZE_KHR]; static cl_device_id deviceId = NULL; - -static const char *vkImage2DShader = - "#version 450\n" - "#extension GL_ARB_separate_shader_objects : enable\n" - "#extension GL_NV_gpu_shader5 : enable\n" - "layout(binding = 0) buffer Params\n" - "{\n" - " uint32_t numImage2DDescriptors;\n" - "};\n" - "layout(binding = 1, " GLSL_FORMAT_STRING - ") uniform " GLSL_TYPE_PREFIX_STRING "image2D image2DList[" STRING( - MAX_2D_IMAGE_DESCRIPTORS) "];\n" - "layout(local_size_x = 32, local_size_y = " - "32) in;\n" - "void main() {\n" - " uvec3 numThreads = gl_NumWorkGroups * " - "gl_WorkGroupSize;\n" - " for (uint32_t image2DIdx = 0; " - "image2DIdx < numImage2DDescriptors; " - "image2DIdx++)" - " {\n" - " ivec2 imageDim = " - "imageSize(image2DList[image2DIdx]);\n" - " uint32_t heightBy2 = imageDim.y / " - "2;\n" - " for (uint32_t row = " - "gl_GlobalInvocationID.y; row < heightBy2; " - "row += numThreads.y)" - " {\n" - " for (uint32_t col = " - "gl_GlobalInvocationID.x; col < imageDim.x; " - "col += numThreads.x)" - " {\n" - " ivec2 coordsA = ivec2(col, " - "row);\n" - " ivec2 coordsB = ivec2(col, " - "imageDim.y - row - 1);\n" - " " GLSL_TYPE_PREFIX_STRING - "vec4 dataA = " - "imageLoad(image2DList[image2DIdx], " - "coordsA);\n" - " " GLSL_TYPE_PREFIX_STRING - "vec4 dataB = " - "imageLoad(image2DList[image2DIdx], " - "coordsB);\n" - " " - "imageStore(image2DList[image2DIdx], " - "coordsA, dataB);\n" - " " - "imageStore(image2DList[image2DIdx], " - "coordsB, dataA);\n" - " }\n" - " }\n" - " }\n" - "}\n"; +size_t max_width = MAX_2D_IMAGE_WIDTH; +size_t max_height = MAX_2D_IMAGE_HEIGHT; const char *kernel_text_numImage_1 = " \ __constant sampler_t smpImg = CLK_NORMALIZED_COORDS_FALSE|CLK_ADDRESS_NONE|CLK_FILTER_NEAREST;\n\ @@ -268,8 +213,8 @@ int run_test_with_two_queue(cl_context &context, cl_command_queue &cmd_queue1, VULKAN_MEMORY_TYPE_PROPERTY_HOST_VISIBLE_COHERENT)); vkParamsDeviceMemory.bindBuffer(vkParamsBuffer); - uint64_t maxImage2DSize = MAX_2D_IMAGE_WIDTH * MAX_2D_IMAGE_HEIGHT - * MAX_2D_IMAGE_ELEMENT_SIZE * 2; + uint64_t maxImage2DSize = + max_width * max_height * MAX_2D_IMAGE_ELEMENT_SIZE * 2; VulkanBuffer vkSrcBuffer(vkDevice, maxImage2DSize); VulkanDeviceMemory vkSrcBufferDeviceMemory( vkDevice, vkSrcBuffer.getSize(), @@ -310,6 +255,12 @@ int run_test_with_two_queue(cl_context &context, cl_command_queue &cmd_queue1, clCl2VkExternalSemaphore = new clExternalSemaphore( vkCl2VkSemaphore, context, vkExternalSemaphoreHandleType, deviceId); + std::vector vkNonDedicatedImage2DListDeviceMemory1; + std::vector vkNonDedicatedImage2DListDeviceMemory2; + std::vector nonDedicatedExternalMemory1; + std::vector nonDedicatedExternalMemory2; + std::vector vkImage2DShader; + for (size_t fIdx = 0; fIdx < vkFormatList.size(); fIdx++) { VulkanFormat vkFormat = vkFormatList[fIdx]; @@ -317,15 +268,13 @@ int run_test_with_two_queue(cl_context &context, cl_command_queue &cmd_queue1, uint32_t elementSize = getVulkanFormatElementSize(vkFormat); ASSERT_LEQ(elementSize, (uint32_t)MAX_2D_IMAGE_ELEMENT_SIZE); log_info("elementSize= %d\n", elementSize); - std::map patternToSubstituteMap; - patternToSubstituteMap[GLSL_FORMAT_STRING] = - getVulkanFormatGLSLFormat(vkFormat); - patternToSubstituteMap[GLSL_TYPE_PREFIX_STRING] = - getVulkanFormatGLSLTypePrefix(vkFormat); - VulkanShaderModule vkImage2DShaderModule( - vkDevice, - prepareVulkanShader(vkImage2DShader, patternToSubstituteMap)); + std::string fileName = "image2D_" + + std::string(getVulkanFormatGLSLFormat(vkFormat)) + ".spv"; + log_info("Load %s file", fileName.c_str()); + vkImage2DShader = readFile(fileName); + VulkanShaderModule vkImage2DShaderModule(vkDevice, vkImage2DShader); + VulkanComputePipeline vkComputePipeline(vkDevice, vkPipelineLayout, vkImage2DShaderModule); @@ -333,13 +282,13 @@ int run_test_with_two_queue(cl_context &context, cl_command_queue &cmd_queue1, { uint32_t width = widthList[wIdx]; log_info("Width: %d\n", width); - ASSERT_LEQ(width, (uint32_t)MAX_2D_IMAGE_WIDTH); + if (width > max_width) continue; region[0] = width; for (size_t hIdx = 0; hIdx < ARRAY_SIZE(heightList); hIdx++) { uint32_t height = heightList[hIdx]; log_info("Height: %d", height); - ASSERT_LEQ(height, (uint32_t)MAX_2D_IMAGE_HEIGHT); + if (height > max_height) continue; region[1] = height; uint32_t numMipLevels = 1; @@ -418,14 +367,6 @@ int run_test_with_two_queue(cl_context &context, cl_command_queue &cmd_queue1, const VulkanMemoryTypeList &memoryTypeList = vkDummyImage2D.getMemoryTypeList(); - std::vector - vkNonDedicatedImage2DListDeviceMemory1; - std::vector - vkNonDedicatedImage2DListDeviceMemory2; - std::vector - nonDedicatedExternalMemory1; - std::vector - nonDedicatedExternalMemory2; for (size_t mtIdx = 0; mtIdx < memoryTypeList.size(); mtIdx++) { @@ -834,6 +775,8 @@ int run_test_with_two_queue(cl_context &context, cl_command_queue &cmd_queue1, } } } + + vkImage2DShader.clear(); } CLEANUP: if (clVk2CLExternalSemaphore) delete clVk2CLExternalSemaphore; @@ -866,8 +809,8 @@ int run_test_with_one_queue(cl_context &context, cl_command_queue &cmd_queue1, VULKAN_MEMORY_TYPE_PROPERTY_HOST_VISIBLE_COHERENT)); vkParamsDeviceMemory.bindBuffer(vkParamsBuffer); - uint64_t maxImage2DSize = MAX_2D_IMAGE_WIDTH * MAX_2D_IMAGE_HEIGHT - * MAX_2D_IMAGE_ELEMENT_SIZE * 2; + uint64_t maxImage2DSize = + max_width * max_height * MAX_2D_IMAGE_ELEMENT_SIZE * 2; VulkanBuffer vkSrcBuffer(vkDevice, maxImage2DSize); VulkanDeviceMemory vkSrcBufferDeviceMemory( vkDevice, vkSrcBuffer.getSize(), @@ -908,6 +851,12 @@ int run_test_with_one_queue(cl_context &context, cl_command_queue &cmd_queue1, clCl2VkExternalSemaphore = new clExternalSemaphore( vkCl2VkSemaphore, context, vkExternalSemaphoreHandleType, deviceId); + std::vector vkNonDedicatedImage2DListDeviceMemory1; + std::vector vkNonDedicatedImage2DListDeviceMemory2; + std::vector nonDedicatedExternalMemory1; + std::vector nonDedicatedExternalMemory2; + std::vector vkImage2DShader; + for (size_t fIdx = 0; fIdx < vkFormatList.size(); fIdx++) { VulkanFormat vkFormat = vkFormatList[fIdx]; @@ -915,15 +864,13 @@ int run_test_with_one_queue(cl_context &context, cl_command_queue &cmd_queue1, uint32_t elementSize = getVulkanFormatElementSize(vkFormat); ASSERT_LEQ(elementSize, (uint32_t)MAX_2D_IMAGE_ELEMENT_SIZE); log_info("elementSize= %d\n", elementSize); - std::map patternToSubstituteMap; - patternToSubstituteMap[GLSL_FORMAT_STRING] = - getVulkanFormatGLSLFormat(vkFormat); - patternToSubstituteMap[GLSL_TYPE_PREFIX_STRING] = - getVulkanFormatGLSLTypePrefix(vkFormat); - VulkanShaderModule vkImage2DShaderModule( - vkDevice, - prepareVulkanShader(vkImage2DShader, patternToSubstituteMap)); + std::string fileName = "image2D_" + + std::string(getVulkanFormatGLSLFormat(vkFormat)) + ".spv"; + log_info("Load %s file", fileName.c_str()); + vkImage2DShader = readFile(fileName); + VulkanShaderModule vkImage2DShaderModule(vkDevice, vkImage2DShader); + VulkanComputePipeline vkComputePipeline(vkDevice, vkPipelineLayout, vkImage2DShaderModule); @@ -931,13 +878,13 @@ int run_test_with_one_queue(cl_context &context, cl_command_queue &cmd_queue1, { uint32_t width = widthList[wIdx]; log_info("Width: %d\n", width); - ASSERT_LEQ(width, (uint32_t)MAX_2D_IMAGE_WIDTH); + if (width > max_width) continue; region[0] = width; for (size_t hIdx = 0; hIdx < ARRAY_SIZE(heightList); hIdx++) { uint32_t height = heightList[hIdx]; log_info("Height: %d\n", height); - ASSERT_LEQ(height, (uint32_t)MAX_2D_IMAGE_HEIGHT); + if (height > max_height) continue; region[1] = height; uint32_t numMipLevels = 1; @@ -1016,14 +963,6 @@ int run_test_with_one_queue(cl_context &context, cl_command_queue &cmd_queue1, const VulkanMemoryTypeList &memoryTypeList = vkDummyImage2D.getMemoryTypeList(); - std::vector - vkNonDedicatedImage2DListDeviceMemory1; - std::vector - vkNonDedicatedImage2DListDeviceMemory2; - std::vector - nonDedicatedExternalMemory1; - std::vector - nonDedicatedExternalMemory2; for (size_t mtIdx = 0; mtIdx < memoryTypeList.size(); mtIdx++) { @@ -1368,6 +1307,7 @@ int run_test_with_one_queue(cl_context &context, cl_command_queue &cmd_queue1, } } } + vkImage2DShader.clear(); } CLEANUP: if (clVk2CLExternalSemaphore) delete clVk2CLExternalSemaphore; @@ -1494,6 +1434,14 @@ int test_image_common(cl_device_id device_, cl_context context_, goto CLEANUP; } deviceId = devices[device_no]; + err = setMaxImageDimensions(deviceId, max_width, max_height); + if (CL_SUCCESS != err) + { + print_error(err, "error setting max image dimensions"); + goto CLEANUP; + } + log_info("Set max_width to %lu and max_height to %lu\n", max_width, + max_height); context = clCreateContextFromType(contextProperties, CL_DEVICE_TYPE_GPU, NULL, NULL, &err); if (CL_SUCCESS != err) diff --git a/test_conformance/vulkan/vulkan_interop_common/opencl_vulkan_wrapper.cpp b/test_conformance/vulkan/vulkan_interop_common/opencl_vulkan_wrapper.cpp index 136818f6..9d9a6601 100644 --- a/test_conformance/vulkan/vulkan_interop_common/opencl_vulkan_wrapper.cpp +++ b/test_conformance/vulkan/vulkan_interop_common/opencl_vulkan_wrapper.cpp @@ -23,6 +23,7 @@ #include #define ASSERT(x) assert((x)) +#define GB(x) ((unsigned long long)(x) << 30) pfnclCreateSemaphoreWithPropertiesKHR clCreateSemaphoreWithPropertiesKHRptr; pfnclEnqueueWaitSemaphoresKHR clEnqueueWaitSemaphoresKHRptr; @@ -31,7 +32,7 @@ pfnclEnqueueAcquireExternalMemObjectsKHR clEnqueueAcquireExternalMemObjectsKHRptr; pfnclEnqueueReleaseExternalMemObjectsKHR clEnqueueReleaseExternalMemObjectsKHRptr; -pfnclReleaseSemaphoreObjectKHR clReleaseSemaphoreObjectKHRptr; +pfnclReleaseSemaphoreKHR clReleaseSemaphoreKHRptr; void init_cl_vk_ext(cl_platform_id opencl_platform) { @@ -51,13 +52,13 @@ void init_cl_vk_ext(cl_platform_id opencl_platform) throw std::runtime_error("Failed to get the function pointer of " "clEnqueueSignalSemaphoresKHRptr!"); } - clReleaseSemaphoreObjectKHRptr = (pfnclReleaseSemaphoreObjectKHR) - clGetExtensionFunctionAddressForPlatform(opencl_platform, - "clReleaseSemaphoreObjectKHR"); - if (NULL == clReleaseSemaphoreObjectKHRptr) + clReleaseSemaphoreKHRptr = + (pfnclReleaseSemaphoreKHR)clGetExtensionFunctionAddressForPlatform( + opencl_platform, "clReleaseSemaphoreKHR"); + if (NULL == clReleaseSemaphoreKHRptr) { throw std::runtime_error("Failed to get the function pointer of " - "clReleaseSemaphoreObjectKHRptr!"); + "clReleaseSemaphoreKHRptr!"); } clCreateSemaphoreWithPropertiesKHRptr = (pfnclCreateSemaphoreWithPropertiesKHR) @@ -70,6 +71,40 @@ void init_cl_vk_ext(cl_platform_id opencl_platform) } } +cl_int setMaxImageDimensions(cl_device_id deviceID, size_t &max_width, + size_t &max_height) +{ + cl_int result = CL_SUCCESS; + cl_ulong val; + size_t paramSize; + + result = clGetDeviceInfo(deviceID, CL_DEVICE_GLOBAL_MEM_SIZE, + sizeof(cl_ulong), &val, ¶mSize); + + if (result != CL_SUCCESS) + { + return result; + } + + if (val < GB(4)) + { + max_width = 256; + max_height = 256; + } + else if (val < GB(8)) + { + max_width = 512; + max_height = 256; + } + else + { + max_width = 1024; + max_height = 512; + } + + return result; +} + cl_int getCLFormatFromVkFormat(VkFormat vkFormat, cl_image_format *clImageFormat) { @@ -798,10 +833,10 @@ clExternalSemaphore::clExternalSemaphore( clExternalSemaphore::~clExternalSemaphore() { - cl_int err = clReleaseSemaphoreObjectKHRptr(m_externalSemaphore); + cl_int err = clReleaseSemaphoreKHRptr(m_externalSemaphore); if (err != CL_SUCCESS) { - throw std::runtime_error("clReleaseSemaphoreObjectKHR failed!"); + throw std::runtime_error("clReleaseSemaphoreKHR failed!"); } } diff --git a/test_conformance/vulkan/vulkan_interop_common/opencl_vulkan_wrapper.hpp b/test_conformance/vulkan/vulkan_interop_common/opencl_vulkan_wrapper.hpp index c1d2a766..d9f8dccb 100644 --- a/test_conformance/vulkan/vulkan_interop_common/opencl_vulkan_wrapper.hpp +++ b/test_conformance/vulkan/vulkan_interop_common/opencl_vulkan_wrapper.hpp @@ -49,7 +49,7 @@ typedef cl_int (*pfnclEnqueueReleaseExternalMemObjectsKHR)( cl_command_queue command_queue, cl_uint num_mem_objects, const cl_mem *mem_objects, cl_uint num_events_in_wait_list, const cl_event *event_wait_list, cl_event *event); -typedef cl_int (*pfnclReleaseSemaphoreObjectKHR)(cl_semaphore_khr sema_object); +typedef cl_int (*pfnclReleaseSemaphoreKHR)(cl_semaphore_khr sema_object); extern pfnclCreateSemaphoreWithPropertiesKHR clCreateSemaphoreWithPropertiesKHRptr; @@ -59,7 +59,7 @@ extern pfnclEnqueueAcquireExternalMemObjectsKHR clEnqueueAcquireExternalMemObjectsKHRptr; extern pfnclEnqueueReleaseExternalMemObjectsKHR clEnqueueReleaseExternalMemObjectsKHRptr; -extern pfnclReleaseSemaphoreObjectKHR clReleaseSemaphoreObjectKHRptr; +extern pfnclReleaseSemaphoreKHR clReleaseSemaphoreKHRptr; cl_int getCLImageInfoFromVkImageInfo(const VkImageCreateInfo *, size_t, cl_image_format *, cl_image_desc *); @@ -69,6 +69,8 @@ cl_int check_external_memory_handle_type( cl_int check_external_semaphore_handle_type( cl_device_id deviceID, cl_external_semaphore_handle_type_khr requiredHandleType); +cl_int setMaxImageDimensions(cl_device_id deviceID, size_t &width, + size_t &height); class clExternalMemory { protected: diff --git a/test_conformance/vulkan/vulkan_interop_common/vulkan_list_map.hpp b/test_conformance/vulkan/vulkan_interop_common/vulkan_list_map.hpp index 831403e1..10a7b221 100644 --- a/test_conformance/vulkan/vulkan_interop_common/vulkan_list_map.hpp +++ b/test_conformance/vulkan/vulkan_interop_common/vulkan_list_map.hpp @@ -335,11 +335,8 @@ const VulkanWrapper & template VulkanWrapper &VulkanList::operator[](size_t idx) { - if (idx < m_wrapperList.size()) - { - // CHECK_LT(idx, m_wrapperList.size()); - return m_wrapperList[idx].get(); - } + // CHECK_LT(idx, m_wrapperList.size()); + return m_wrapperList[idx].get(); } template diff --git a/test_conformance/vulkan/vulkan_interop_common/vulkan_utility.cpp b/test_conformance/vulkan/vulkan_interop_common/vulkan_utility.cpp index 81e12621..4e6118b1 100644 --- a/test_conformance/vulkan/vulkan_interop_common/vulkan_utility.cpp +++ b/test_conformance/vulkan/vulkan_interop_common/vulkan_utility.cpp @@ -18,6 +18,7 @@ #include "vulkan_wrapper.hpp" #include #include +#include #include #include #include @@ -541,59 +542,6 @@ const char *getVulkanFormatGLSLFormat(VulkanFormat format) return (const char *)size_t(0); } -const char *getVulkanFormatGLSLTypePrefix(VulkanFormat format) -{ - switch (format) - { - case VULKAN_FORMAT_R8_UINT: - case VULKAN_FORMAT_R8G8_UINT: - case VULKAN_FORMAT_R8G8B8A8_UINT: - case VULKAN_FORMAT_R16_UINT: - case VULKAN_FORMAT_R16G16_UINT: - case VULKAN_FORMAT_R16G16B16A16_UINT: - case VULKAN_FORMAT_R32_UINT: - case VULKAN_FORMAT_R32G32_UINT: - case VULKAN_FORMAT_R32G32B32A32_UINT: return "u"; - - case VULKAN_FORMAT_R8_SINT: - case VULKAN_FORMAT_R8G8_SINT: - case VULKAN_FORMAT_R8G8B8A8_SINT: - case VULKAN_FORMAT_R16_SINT: - case VULKAN_FORMAT_R16G16_SINT: - case VULKAN_FORMAT_R16G16B16A16_SINT: - case VULKAN_FORMAT_R32_SINT: - case VULKAN_FORMAT_R32G32_SINT: - case VULKAN_FORMAT_R32G32B32A32_SINT: return "i"; - - case VULKAN_FORMAT_R32_SFLOAT: - case VULKAN_FORMAT_R32G32_SFLOAT: - case VULKAN_FORMAT_R32G32B32A32_SFLOAT: return ""; - - default: ASSERT(0); std::cout << "Unknown format"; - } - - return ""; -} - -std::string prepareVulkanShader( - std::string shaderCode, - const std::map &patternToSubstituteMap) -{ - for (std::map::const_iterator psIt = - patternToSubstituteMap.begin(); - psIt != patternToSubstituteMap.end(); ++psIt) - { - std::string::size_type pos = 0u; - while ((pos = shaderCode.find(psIt->first, pos)) != std::string::npos) - { - shaderCode.replace(pos, psIt->first.length(), psIt->second); - pos += psIt->second.length(); - } - } - - return shaderCode; -} - std::ostream &operator<<(std::ostream &os, VulkanMemoryTypeProperty memoryTypeProperty) { @@ -691,3 +639,54 @@ std::ostream &operator<<(std::ostream &os, VulkanFormat format) return os; } + +static char *findFilePath(const std::string filename) +{ + const char *searchPath[] = { + "./", // Same dir + "./shaders/", // In shaders folder in same dir + "../test_conformance/vulkan/shaders/" // In src folder + }; + for (unsigned int i = 0; i < sizeof(searchPath) / sizeof(char *); ++i) + { + std::string path(searchPath[i]); + + path.append(filename); + FILE *fp; + fp = fopen(path.c_str(), "rb"); + + if (fp != NULL) + { + fclose(fp); + // File found + char *file_path = (char *)(malloc(path.length() + 1)); + strncpy(file_path, path.c_str(), path.length() + 1); + return file_path; + } + if (fp) + { + fclose(fp); + } + } + // File not found + return 0; +} + +std::vector readFile(const std::string &filename) +{ + char *file_path = findFilePath(filename); + + std::ifstream file(file_path, std::ios::ate | std::ios::binary); + + if (!file.is_open()) + { + throw std::runtime_error("failed to open shader spv file!\n"); + } + size_t fileSize = (size_t)file.tellg(); + std::vector buffer(fileSize); + file.seekg(0); + file.read(buffer.data(), fileSize); + file.close(); + printf("filesize is %d", fileSize); + return buffer; +} diff --git a/test_conformance/vulkan/vulkan_interop_common/vulkan_utility.hpp b/test_conformance/vulkan/vulkan_interop_common/vulkan_utility.hpp index 7022fd5a..04f5a594 100644 --- a/test_conformance/vulkan/vulkan_interop_common/vulkan_utility.hpp +++ b/test_conformance/vulkan/vulkan_interop_common/vulkan_utility.hpp @@ -66,4 +66,5 @@ operator<<(std::ostream& os, VulkanExternalSemaphoreHandleType externalSemaphoreHandleType); std::ostream& operator<<(std::ostream& os, VulkanFormat format); +std::vector readFile(const std::string& filename); #endif // _vulkan_utility_hpp_ diff --git a/test_conformance/vulkan/vulkan_interop_common/vulkan_wrapper.cpp b/test_conformance/vulkan/vulkan_interop_common/vulkan_wrapper.cpp index c044e009..e5d3a271 100644 --- a/test_conformance/vulkan/vulkan_interop_common/vulkan_wrapper.cpp +++ b/test_conformance/vulkan/vulkan_interop_common/vulkan_wrapper.cpp @@ -201,7 +201,8 @@ VulkanInstance::VulkanInstance(): m_vkInstance(VK_NULL_HANDLE) if (physicalDeviceCount == uint32_t(0)) { - throw std::runtime_error("failed to find GPUs with Vulkan support!"); + std::cout << "failed to find GPUs with Vulkan support!\n"; + return; } std::vector vkPhysicalDeviceList(physicalDeviceCount, @@ -846,23 +847,18 @@ VulkanShaderModule::VulkanShaderModule(const VulkanShaderModule &shaderModule) {} VulkanShaderModule::VulkanShaderModule(const VulkanDevice &device, - const std::string &code) + const std::vector &code) : m_device(device) { - std::string paddedCode = code; - while (paddedCode.size() % 4) - { - paddedCode += " "; - } VkShaderModuleCreateInfo vkShaderModuleCreateInfo = {}; vkShaderModuleCreateInfo.sType = VK_STRUCTURE_TYPE_SHADER_MODULE_CREATE_INFO; vkShaderModuleCreateInfo.pNext = NULL; vkShaderModuleCreateInfo.flags = 0; - vkShaderModuleCreateInfo.codeSize = paddedCode.size(); + vkShaderModuleCreateInfo.codeSize = code.size(); vkShaderModuleCreateInfo.pCode = - (const uint32_t *)(void *)paddedCode.c_str(); + reinterpret_cast(code.data()); vkCreateShaderModule(m_device, &vkShaderModuleCreateInfo, NULL, &m_vkShaderModule); diff --git a/test_conformance/vulkan/vulkan_interop_common/vulkan_wrapper.hpp b/test_conformance/vulkan/vulkan_interop_common/vulkan_wrapper.hpp index 1f68a92b..37925ee4 100644 --- a/test_conformance/vulkan/vulkan_interop_common/vulkan_wrapper.hpp +++ b/test_conformance/vulkan/vulkan_interop_common/vulkan_wrapper.hpp @@ -240,7 +240,8 @@ protected: VulkanShaderModule(const VulkanShaderModule &shaderModule); public: - VulkanShaderModule(const VulkanDevice &device, const std::string &code); + VulkanShaderModule(const VulkanDevice &device, + const std::vector &code); virtual ~VulkanShaderModule(); operator VkShaderModule() const; };