diff --git a/presubmit.sh b/presubmit.sh index 10354abf14..23e1e9a3b7 100755 --- a/presubmit.sh +++ b/presubmit.sh @@ -62,7 +62,7 @@ cmake .. -G Ninja \ -DCMAKE_BUILD_TYPE=Release \ -DCMAKE_TOOLCHAIN_FILE=${TOOLCHAIN_FILE} \ -DOPENCL_ICD_LOADER_HEADERS_DIR=${TOP}/OpenCL-Headers/ -cmake --build . -j2 +cmake --build . --parallel #Vulkan Loader cd ${TOP} @@ -78,7 +78,7 @@ cmake .. -G Ninja \ -DBUILD_WSI_XCB_SUPPORT=OFF \ -DBUILD_WSI_WAYLAND_SUPPORT=OFF \ -C helper.cmake .. -cmake --build . -j2 +cmake --build . --parallel # Build CTS cd ${TOP} @@ -105,4 +105,4 @@ cmake .. -G Ninja \ -DVULKAN_IS_SUPPORTED=${BUILD_VULKAN_TEST} \ -DVULKAN_INCLUDE_DIR=${TOP}/Vulkan-Headers/include/ \ -DVULKAN_LIB_DIR=${TOP}/Vulkan-Loader/build/loader/ -cmake --build . -j3 +cmake --build . --parallel diff --git a/test_conformance/basic/CMakeLists.txt b/test_conformance/basic/CMakeLists.txt index 684a7d1d4b..bf1f3bd63a 100644 --- a/test_conformance/basic/CMakeLists.txt +++ b/test_conformance/basic/CMakeLists.txt @@ -26,7 +26,6 @@ set(${MODULE_NAME}_SOURCES test_arrayreadwrite.cpp test_arraycopy.cpp test_imagearraycopy.cpp - test_imagearraycopy3d.cpp test_imagecopy.cpp test_imagerandomcopy.cpp test_arrayimagecopy.cpp diff --git a/test_conformance/basic/test_if.cpp b/test_conformance/basic/test_if.cpp index c92ec32218..f2a8fa8299 100644 --- a/test_conformance/basic/test_if.cpp +++ b/test_conformance/basic/test_if.cpp @@ -1,6 +1,6 @@ // // Copyright (c) 2017 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 @@ -21,146 +21,119 @@ #include #include +#include +#include #include "procs.h" -const char *conditional_kernel_code = -"__kernel void test_if(__global int *src, __global int *dst)\n" -"{\n" -" int tid = get_global_id(0);\n" -"\n" -" if (src[tid] == 0)\n" -" dst[tid] = 0x12345678;\n" -" else if (src[tid] == 1)\n" -" dst[tid] = 0x23456781;\n" -" else if (src[tid] == 2)\n" -" dst[tid] = 0x34567812;\n" -" else if (src[tid] == 3)\n" -" dst[tid] = 0x45678123;\n" -" else if (src[tid] == 4)\n" -" dst[tid] = 0x56781234;\n" -" else if (src[tid] == 5)\n" -" dst[tid] = 0x67812345;\n" -" else if (src[tid] == 6)\n" -" dst[tid] = 0x78123456;\n" -" else if (src[tid] == 7)\n" -" dst[tid] = 0x81234567;\n" -" else\n" -" dst[tid] = 0x7FFFFFFF;\n" -"\n" -"}\n"; - -const int results[] = { - 0x12345678, - 0x23456781, - 0x34567812, - 0x45678123, - 0x56781234, - 0x67812345, - 0x78123456, - 0x81234567, -}; - -int -verify_if(int *inptr, int *outptr, int n) +namespace { +const char *conditional_kernel_code = R"( +__kernel void test_if(__global int *src, __global int *dst) { - int r, i; + int tid = get_global_id(0); + + if (src[tid] == 0) + dst[tid] = 0x12345678; + else if (src[tid] == 1) + dst[tid] = 0x23456781; + else if (src[tid] == 2) + dst[tid] = 0x34567812; + else if (src[tid] == 3) + dst[tid] = 0x45678123; + else if (src[tid] == 4) + dst[tid] = 0x56781234; + else if (src[tid] == 5) + dst[tid] = 0x67812345; + else if (src[tid] == 6) + dst[tid] = 0x78123456; + else if (src[tid] == 7) + dst[tid] = 0x81234567; + else + dst[tid] = 0x7FFFFFFF; +} +)"; - for (i=0; i input, std::vector output) +{ + const cl_int results[] = { + 0x12345678, 0x23456781, 0x34567812, 0x45678123, + 0x56781234, 0x67812345, 0x78123456, 0x81234567, + }; + + auto predicate = [&results](cl_int a, cl_int b) { + if (a <= 7) + return b == results[a]; else - r = 0x7FFFFFFF; + return b == 0x7FFFFFFF; + }; - if (r != outptr[i]) - { - log_error("IF test failed\n"); - return -1; - } + if (!std::equal(input.begin(), input.end(), output.begin(), predicate)) + { + log_error("IF test failed\n"); + return -1; } log_info("IF test passed\n"); return 0; } -int test_if(cl_device_id device, cl_context context, cl_command_queue queue, int num_elements) +void generate_random_inputs(std::vector &v) { - cl_mem streams[2]; - cl_int *input_ptr, *output_ptr; - cl_program program; - cl_kernel kernel; - size_t threads[1]; - int err, i; - MTdata d = init_genrand( gRandomSeed ); + RandomSeed seed(gRandomSeed); + + auto random_generator = [&seed]() { + return static_cast(get_random_float(0, 32, seed)); + }; + + std::generate(v.begin(), v.end(), random_generator); +} +} +int test_if(cl_device_id device, cl_context context, cl_command_queue queue, + int num_elements) +{ + clMemWrapper streams[2]; + clProgramWrapper program; + clKernelWrapper kernel; + + int err; size_t length = sizeof(cl_int) * num_elements; - input_ptr = (cl_int*)malloc(length); - output_ptr = (cl_int*)malloc(length); - streams[0] = clCreateBuffer(context, CL_MEM_READ_WRITE, length, NULL, NULL); - if (!streams[0]) - { - log_error("clCreateBuffer failed\n"); - return -1; - } - streams[1] = clCreateBuffer(context, CL_MEM_READ_WRITE, length, NULL, NULL); - if (!streams[1]) - { - log_error("clCreateBuffer failed\n"); - return -1; - } + std::vector input(num_elements); + std::vector output(num_elements); - for (i=0; i #include #include #include #include +#include #include "procs.h" -int test_imagearraycopy_single_format(cl_device_id device, cl_context context, cl_command_queue queue, cl_image_format *format) +int test_imagearraycopy_single_format(cl_device_id device, cl_context context, + cl_command_queue queue, + cl_mem_flags flags, + cl_mem_object_type image_type, + const cl_image_format *format) { - cl_uchar *imgptr, *bufptr; - clMemWrapper image, buffer; - int img_width = 512; - int img_height = 512; - size_t elem_size; - size_t buffer_size; - int i; - cl_int err; - MTdata d; - cl_event copyevent; - - log_info("Testing %s %s\n", GetChannelOrderName(format->image_channel_order), GetChannelTypeName(format->image_channel_data_type)); - - image = create_image_2d(context, CL_MEM_READ_WRITE, format, img_width, - img_height, 0, NULL, &err); - test_error(err, "create_image_2d failed"); - - err = clGetImageInfo(image, CL_IMAGE_ELEMENT_SIZE, sizeof(size_t), &elem_size, NULL); - test_error(err, "clGetImageInfo failed"); - - buffer_size = sizeof(cl_uchar) * elem_size * img_width * img_height; - - buffer = clCreateBuffer(context, CL_MEM_READ_WRITE, buffer_size, NULL, &err); - test_error(err, "clCreateBuffer failed"); - - d = init_genrand( gRandomSeed ); - imgptr = (cl_uchar*)malloc(buffer_size); - for (i=0; i<(int)buffer_size; i++) { - imgptr[i] = (cl_uchar)genrand_int32(d); - } - free_mtdata(d); d = NULL; - - size_t origin[3]={0,0,0}, region[3]={img_width,img_height,1}; - err = clEnqueueWriteImage( queue, image, CL_TRUE, origin, region, 0, 0, imgptr, 0, NULL, NULL ); - test_error(err, "clEnqueueWriteBuffer failed"); - - err = clEnqueueCopyImageToBuffer( queue, image, buffer, origin, region, 0, 0, NULL, ©event ); - test_error(err, "clEnqueueCopyImageToBuffer failed"); - - bufptr = (cl_uchar*)malloc(buffer_size); - - err = clEnqueueReadBuffer( queue, buffer, CL_TRUE, 0, buffer_size, bufptr, 1, ©event, NULL); - test_error(err, "clEnqueueReadBuffer failed"); - - err = clReleaseEvent(copyevent); - test_error(err, "clReleaseEvent failed"); - - if (memcmp(imgptr, bufptr, buffer_size) != 0) { - log_error( "ERROR: Results did not validate!\n" ); - unsigned char * inchar = (unsigned char*)imgptr; - unsigned char * outchar = (unsigned char*)bufptr; - int failuresPrinted = 0; - int i; - for (i=0; i< (int)buffer_size; i+=(int)elem_size) { - int failed = 0; - int j; - for (j=0; j<(int)elem_size; j++) - if (inchar[i+j] != outchar[i+j]) - failed = 1; - char values[4096]; - values[0] = 0; - if (failed) { - sprintf(values + strlen(values), "%d(0x%x) -> expected [", i, i); - int j; - for (j=0; j<(int)elem_size; j++) - sprintf(values + strlen( values), "0x%02x ", inchar[i+j]); - sprintf(values + strlen(values), "] != actual ["); - for (j=0; j<(int)elem_size; j++) - sprintf(values + strlen( values), "0x%02x ", outchar[i+j]); - sprintf(values + strlen(values), "]"); - log_error("%s\n", values); - failuresPrinted++; - } - if (failuresPrinted > 5) { - log_error("Not printing further failures...\n"); - break; + clMemWrapper buffer, image; + const int img_width = 512; + const int img_height = 512; + const int img_depth = (image_type == CL_MEM_OBJECT_IMAGE3D) ? 32 : 1; + size_t elem_size; + size_t buffer_size; + cl_int err; + cl_event copyevent; + + log_info("Testing %s %s\n", + GetChannelOrderName(format->image_channel_order), + GetChannelTypeName(format->image_channel_data_type)); + + if (CL_MEM_OBJECT_IMAGE2D == image_type) + { + image = create_image_2d(context, flags, format, img_width, img_height, + 0, nullptr, &err); + } + else + { + image = create_image_3d(context, flags, format, img_width, img_height, + img_depth, 0, 0, nullptr, &err); + } + test_error(err, "create_image_xd failed"); + + err = clGetImageInfo(image, CL_IMAGE_ELEMENT_SIZE, sizeof(size_t), + &elem_size, nullptr); + test_error(err, "clGetImageInfo failed"); + + buffer_size = + sizeof(cl_uchar) * elem_size * img_width * img_height * img_depth; + + buffer = + clCreateBuffer(context, CL_MEM_READ_WRITE, buffer_size, nullptr, &err); + test_error(err, "clCreateBuffer failed"); + + + RandomSeed seed(gRandomSeed); + cl_uchar *imgptr = + static_cast(create_random_data(kUChar, seed, buffer_size)); + + const size_t origin[3] = { 0, 0, 0 }, + region[3] = { img_width, img_height, img_depth }; + err = clEnqueueWriteImage(queue, image, CL_TRUE, origin, region, 0, 0, + imgptr, 0, nullptr, nullptr); + test_error(err, "clEnqueueWriteImage failed"); + + err = clEnqueueCopyImageToBuffer(queue, image, buffer, origin, region, 0, 0, + nullptr, ©event); + test_error(err, "clEnqueueCopyImageToBuffer failed"); + + cl_uchar *bufptr = static_cast(malloc(buffer_size)); + + err = clEnqueueReadBuffer(queue, buffer, CL_TRUE, 0, buffer_size, bufptr, 1, + ©event, nullptr); + test_error(err, "clEnqueueReadBuffer failed"); + + err = clReleaseEvent(copyevent); + test_error(err, "clReleaseEvent failed"); + + if (memcmp(imgptr, bufptr, buffer_size) != 0) + { + log_error("ERROR: Results did not validate!\n"); + auto inchar = static_cast(imgptr); + auto outchar = static_cast(bufptr); + int failuresPrinted = 0; + for (size_t i = 0; i < buffer_size; i += elem_size) + { + if (memcmp(&inchar[i], &outchar[i], elem_size) != 0) + { + log_error("%d(0x%x) -> expected [", i, i); + for (size_t j = 0; j < elem_size; j++) + log_error("0x%02x ", inchar[i + j]); + log_error("] != actual ["); + for (size_t j = 0; j < elem_size; j++) + log_error("0x%02x ", outchar[i + j]); + log_error("]\n"); + failuresPrinted++; + } + if (failuresPrinted > 5) + { + log_error("Not printing further failures...\n"); + break; + } } + err = -1; } - err = -1; - } - free(imgptr); - free(bufptr); + free(imgptr); + free(bufptr); - if (err) - log_error("IMAGE to ARRAY copy test failed for image_channel_order=0x%lx and image_channel_data_type=0x%lx\n", - (unsigned long)format->image_channel_order, (unsigned long)format->image_channel_data_type); + if (err) + log_error( + "IMAGE to ARRAY copy test failed for image_channel_order=0x%lx and " + "image_channel_data_type=0x%lx\n", + static_cast(format->image_channel_order), + static_cast(format->image_channel_data_type)); - return err; + return err; } -int test_imagearraycopy(cl_device_id device, cl_context context, cl_command_queue queue, int num_elements) +int test_imagearraycommon(cl_device_id device, cl_context context, + cl_command_queue queue, cl_mem_flags flags, + cl_mem_object_type image_type) { - cl_int err; - cl_image_format *formats; - cl_uint num_formats; - cl_uint i; + cl_int err; + cl_uint num_formats; + + err = clGetSupportedImageFormats(context, flags, image_type, 0, nullptr, + &num_formats); + test_error(err, "clGetSupportedImageFormats failed"); - PASSIVE_REQUIRE_IMAGE_SUPPORT( device ) + std::vector formats(num_formats); - err = clGetSupportedImageFormats(context, CL_MEM_READ_WRITE, CL_MEM_OBJECT_IMAGE2D, 0, NULL, &num_formats); - test_error(err, "clGetSupportedImageFormats failed"); + err = clGetSupportedImageFormats(context, flags, image_type, num_formats, + formats.data(), nullptr); + test_error(err, "clGetSupportedImageFormats failed"); - formats = (cl_image_format *)malloc(num_formats * sizeof(cl_image_format)); + for (const auto &format : formats) + { + err |= test_imagearraycopy_single_format(device, context, queue, flags, + image_type, &format); + } - err = clGetSupportedImageFormats(context, CL_MEM_READ_WRITE, CL_MEM_OBJECT_IMAGE2D, num_formats, formats, NULL); - test_error(err, "clGetSupportedImageFormats failed"); + if (err) + log_error("ARRAY to IMAGE%s copy test failed\n", + convert_image_type_to_string(image_type)); + else + log_info("ARRAY to IMAGE%s copy test passed\n", + convert_image_type_to_string(image_type)); - for (i = 0; i < num_formats; i++) { - err |= test_imagearraycopy_single_format(device, context, queue, &formats[i]); - } + return err; +} - free(formats); - if (err) - log_error("IMAGE to ARRAY copy test failed\n"); - else - log_info("IMAGE to ARRAY copy test passed\n"); +int test_imagearraycopy(cl_device_id device, cl_context context, + cl_command_queue queue, int num_elements) +{ + PASSIVE_REQUIRE_IMAGE_SUPPORT(device) - return err; + return test_imagearraycommon(device, context, queue, CL_MEM_READ_WRITE, + CL_MEM_OBJECT_IMAGE2D); } + + +int test_imagearraycopy3d(cl_device_id device, cl_context context, + cl_command_queue queue, int num_elements) +{ + PASSIVE_REQUIRE_3D_IMAGE_SUPPORT(device) + + return test_imagearraycommon(device, context, queue, CL_MEM_READ_ONLY, + CL_MEM_OBJECT_IMAGE3D); +} \ No newline at end of file diff --git a/test_conformance/basic/test_imagearraycopy3d.cpp b/test_conformance/basic/test_imagearraycopy3d.cpp deleted file mode 100644 index 60b8a58419..0000000000 --- a/test_conformance/basic/test_imagearraycopy3d.cpp +++ /dev/null @@ -1,147 +0,0 @@ -// -// Copyright (c) 2017 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 "harness/compat.h" - -#include -#include -#include -#include -#include - -#include "procs.h" - -int test_imagearraycopy3d_single_format(cl_device_id device, cl_context context, cl_command_queue queue, cl_image_format *format) -{ - cl_uchar *imgptr, *bufptr; - clMemWrapper image, buffer; - int img_width = 128; - int img_height = 128; - int img_depth = 32; - size_t elem_size; - size_t buffer_size; - int i; - cl_int err; - MTdata d; - - log_info("Testing %s %s\n", GetChannelOrderName(format->image_channel_order), GetChannelTypeName(format->image_channel_data_type)); - - image = create_image_3d(context, CL_MEM_READ_ONLY, format, img_width, - img_height, img_depth, 0, 0, NULL, &err); - test_error(err, "create_image_3d failed"); - - err = clGetImageInfo(image, CL_IMAGE_ELEMENT_SIZE, sizeof(size_t), &elem_size, NULL); - test_error(err, "clGetImageInfo failed"); - - buffer_size = sizeof(cl_uchar) * elem_size * img_width * img_height * img_depth; - - buffer = clCreateBuffer(context, CL_MEM_READ_WRITE, buffer_size, NULL, &err); - test_error(err, "clCreateBuffer failed"); - - d = init_genrand( gRandomSeed ); - imgptr = (cl_uchar*)malloc(buffer_size); - for (i=0; i<(int)buffer_size; i++) { - imgptr[i] = (cl_uchar)genrand_int32(d); - } - free_mtdata(d); d = NULL; - - size_t origin[3]={0,0,0}, region[3]={img_width,img_height,img_depth}; - err = clEnqueueWriteImage( queue, image, CL_TRUE, origin, region, 0, 0, imgptr, 0, NULL, NULL ); - test_error(err, "clEnqueueWriteBuffer failed"); - - err = clEnqueueCopyImageToBuffer( queue, image, buffer, origin, region, 0, 0, NULL, NULL ); - test_error(err, "clEnqueueCopyImageToBuffer failed"); - - bufptr = (cl_uchar*)malloc(buffer_size); - - err = clEnqueueReadBuffer( queue, buffer, CL_TRUE, 0, buffer_size, bufptr, 0, NULL, NULL); - test_error(err, "clEnqueueReadBuffer failed"); - - if (memcmp(imgptr, bufptr, buffer_size) != 0) { - log_error( "ERROR: Results did not validate!\n" ); - unsigned char * inchar = (unsigned char*)imgptr; - unsigned char * outchar = (unsigned char*)bufptr; - int failuresPrinted = 0; - int i; - for (i=0; i< (int)buffer_size; i+=(int)elem_size) { - int failed = 0; - int j; - for (j=0; j<(int)elem_size; j++) - if (inchar[i+j] != outchar[i+j]) - failed = 1; - char values[4096]; - values[0] = 0; - if (failed) { - sprintf(values + strlen(values), "%d(0x%x) -> expected [", i, i); - int j; - for (j=0; j<(int)elem_size; j++) - sprintf(values + strlen( values), "0x%02x ", inchar[i+j]); - sprintf(values + strlen(values), "] != actual ["); - for (j=0; j<(int)elem_size; j++) - sprintf(values + strlen( values), "0x%02x ", outchar[i+j]); - sprintf(values + strlen(values), "]"); - log_error("%s\n", values); - failuresPrinted++; - } - if (failuresPrinted > 5) { - log_error("Not printing further failures...\n"); - break; - } - } - err = -1; - } - - free(imgptr); - free(bufptr); - - if (err) - log_error("IMAGE3D to ARRAY copy test failed for image_channel_order=0x%lx and image_channel_data_type=0x%lx\n", - (unsigned long)format->image_channel_order, (unsigned long)format->image_channel_data_type); - - return err; -} - -int test_imagearraycopy3d(cl_device_id device, cl_context context, cl_command_queue queue, int num_elements) -{ - cl_int err; - cl_image_format *formats; - cl_uint num_formats; - cl_uint i; - - PASSIVE_REQUIRE_3D_IMAGE_SUPPORT( device ) - - err = clGetSupportedImageFormats( - context, CL_MEM_READ_ONLY, CL_MEM_OBJECT_IMAGE3D, 0, NULL, &num_formats); - test_error(err, "clGetSupportedImageFormats failed"); - - formats = (cl_image_format *)malloc(num_formats * sizeof(cl_image_format)); - - err = clGetSupportedImageFormats(context, CL_MEM_READ_ONLY, - CL_MEM_OBJECT_IMAGE3D, num_formats, formats, - NULL); - test_error(err, "clGetSupportedImageFormats failed"); - - for (i = 0; i < num_formats; i++) { - err |= test_imagearraycopy3d_single_format(device, context, queue, &formats[i]); - } - - free(formats); - if (err) - log_error("IMAGE3D to ARRAY copy test failed\n"); - else - log_info("IMAGE3D to ARRAY copy test passed\n"); - - return err; -} diff --git a/test_conformance/basic/test_imagedim.cpp b/test_conformance/basic/test_imagedim.cpp index 008c88b6af..f979aa8bb1 100644 --- a/test_conformance/basic/test_imagedim.cpp +++ b/test_conformance/basic/test_imagedim.cpp @@ -1,6 +1,6 @@ // // Copyright (c) 2017 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 @@ -21,504 +21,277 @@ #include #include +#include +#include #include "procs.h" -static const char *image_dim_kernel_code = -"\n" -"__kernel void test_image_dim(read_only image2d_t srcimg, write_only image2d_t dstimg, sampler_t sampler)\n" -"{\n" -" int tid_x = get_global_id(0);\n" -" int tid_y = get_global_id(1);\n" -" float4 color;\n" -"\n" -" color = read_imagef(srcimg, sampler, (int2)(tid_x, tid_y));\n" -" write_imagef(dstimg, (int2)(tid_x, tid_y), color);\n" -"\n" -"}\n"; - - -static unsigned char *generate_8888_image(size_t w, size_t h, MTdata d) +namespace { +const char *image_dim_kernel_code = R"( +__kernel void test_image_dim(read_only image2d_t srcimg, write_only image2d_t dstimg, sampler_t sampler) { - unsigned char *ptr = new unsigned char[4 * w * h]; - size_t i; - - for (i = 0; i < w * h * 4; i++) - { - ptr[i] = (unsigned char)genrand_int32(d); - } + int tid_x = get_global_id(0); + int tid_y = get_global_id(1); + float4 color; - return ptr; + color = read_imagef(srcimg, sampler, (int2)(tid_x, tid_y)); + write_imagef(dstimg, (int2)(tid_x, tid_y), color); } +)"; -static int verify_8888_image(unsigned char *image, unsigned char *outptr, - size_t w, size_t h) +void generate_random_inputs(std::vector &v) { - size_t i; + RandomSeed seed(gRandomSeed); - for (i = 0; i < w * h; i++) - { - if (outptr[i] != image[i]) - return -1; - } + auto random_generator = [&seed]() { return genrand_int32(seed); }; - return 0; + std::generate(v.begin(), v.end(), random_generator); } - -int -test_imagedim_pow2(cl_device_id device, cl_context context, cl_command_queue queue, int n_elems) +int get_max_image_dimensions(cl_device_id device, size_t &max_img_width, + size_t &max_img_height) { - cl_mem streams[2]; - cl_image_format img_format; - unsigned char *input_ptr, *output_ptr; - cl_program program; - cl_kernel kernel; - size_t threads[2]; + int err = 0; + cl_ulong max_mem_size; - size_t img_width, max_img_width; - size_t img_height, max_img_height; - size_t max_img_dim; - int i, j, i2, j2, err = 0; size_t max_image2d_width, max_image2d_height; - int total_errors = 0; - MTdata d; - - PASSIVE_REQUIRE_IMAGE_SUPPORT( device ) - err = create_single_kernel_helper( context, &program, &kernel, 1, &image_dim_kernel_code, "test_image_dim" ); - if (err) - { - log_error("create_program_and_kernel_with_sources failed\n"); - return -1; - } - - err = clGetDeviceInfo(device, CL_DEVICE_GLOBAL_MEM_SIZE,sizeof(max_mem_size), &max_mem_size, NULL); - if (err) - { - log_error("clGetDeviceInfo for CL_DEVICE_GLOBAL_MEM_SIZE failed (%d)\n", err); - return -1; - } - err = clGetDeviceInfo(device, CL_DEVICE_IMAGE2D_MAX_WIDTH, sizeof(max_image2d_width), &max_image2d_width, NULL); - if (err) - { - log_error("clGetDeviceInfo for CL_DEVICE_IMAGE2D_MAX_WIDTH failed (%d)\n", err); - return -1; - } - err = clGetDeviceInfo(device, CL_DEVICE_IMAGE2D_MAX_HEIGHT, sizeof(max_image2d_width), &max_image2d_height, NULL); - if (err) - { - log_error("clGetDeviceInfo for CL_DEVICE_IMAGE2D_MAX_HEIGHT failed (%d)\n", err); - return -1; - } - log_info("Device reported max image sizes of %lu x %lu, and max mem size of %gMB.\n", - max_image2d_width, max_image2d_height, max_mem_size/(1024.0*1024.0)); - - if (max_mem_size > (cl_ulong)SIZE_MAX) { - max_mem_size = (cl_ulong)SIZE_MAX; - } - - cl_sampler sampler = clCreateSampler(context, CL_FALSE, CL_ADDRESS_CLAMP_TO_EDGE, CL_FILTER_NEAREST, &err); - test_error(err, "clCreateSampler failed"); - - max_img_width = max_image2d_width; - max_img_height = max_image2d_height; - - // determine max image dim we can allocate - assume RGBA image, 4 bytes per pixel, - // and we want to consume 1/4 of global memory (this is the minimum required to be - // supported by the spec) + err = clGetDeviceInfo(device, CL_DEVICE_GLOBAL_MEM_SIZE, + sizeof(max_mem_size), &max_mem_size, nullptr); + test_error(err, "clGetDeviceInfo for CL_DEVICE_GLOBAL_MEM_SIZE failed"); + err = + clGetDeviceInfo(device, CL_DEVICE_IMAGE2D_MAX_WIDTH, + sizeof(max_image2d_width), &max_image2d_width, nullptr); + test_error(err, "clGetDeviceInfo for CL_DEVICE_IMAGE2D_MAX_WIDTH failed"); + err = clGetDeviceInfo(device, CL_DEVICE_IMAGE2D_MAX_HEIGHT, + sizeof(max_image2d_width), &max_image2d_height, + nullptr); + test_error(err, "clGetDeviceInfo for CL_DEVICE_IMAGE2D_MAX_HEIGHT failed"); + + log_info("Device reported max image sizes of %lu x %lu, and max mem size " + "of %gMB.\n", + max_image2d_width, max_image2d_height, + max_mem_size / (1024.0 * 1024.0)); + + + max_mem_size = std::min(max_mem_size, (cl_ulong)SIZE_MAX); + + // determine max image dim we can allocate - assume RGBA image, 4 bytes per + // pixel, and we want to consume 1/4 of global memory (this is the minimum + // required to be supported by the spec) max_mem_size /= 4; // use 1/4 max_mem_size /= 4; // 4 bytes per pixel - max_img_dim = (size_t)sqrt((double)max_mem_size); + + size_t max_img_dim = + static_cast(sqrt(static_cast(max_mem_size))); // convert to a power of 2 { - unsigned int n = (unsigned int)max_img_dim; - unsigned int m = 0x80000000; + unsigned int n = static_cast(max_img_dim); + unsigned int m = 0x80000000; // round-down to the nearest power of 2 - while (m > n) - m >>= 1; + while (m > n) m >>= 1; max_img_dim = m; } - if (max_img_width > max_img_dim) - max_img_width = max_img_dim; - if (max_img_height > max_img_dim) - max_img_height = max_img_dim; + max_img_width = std::min(max_image2d_width, max_img_dim); + max_img_height = std::min(max_image2d_height, max_img_dim); - log_info("Adjusted maximum image size to test is %d x %d, which is a max mem size of %gMB.\n", - max_img_width, max_img_height, (max_img_width*max_img_height*4)/(1024.0*1024.0)); + log_info("Adjusted maximum image size to test is %d x %d, which is a max " + "mem size of %gMB.\n", + max_img_width, max_img_height, + (max_img_width * max_img_height * 4) / (1024.0 * 1024.0)); + return err; +} - d = init_genrand( gRandomSeed ); - input_ptr = generate_8888_image(max_img_width, max_img_height, d); +int test_imagedim_common(cl_context context, cl_command_queue queue, + cl_kernel kernel, size_t *local_threads, + size_t img_width, size_t img_height) +{ - output_ptr = new unsigned char[4 * max_img_width * max_img_height]; + int err; + int total_errors = 0; - // test power of 2 width, height starting at 1 to 4K - for (i = 1, i2 = 0; i <= max_img_height; i <<= 1, i2++) - { - img_height = (1 << i2); - for (j = 1, j2 = 0; j <= max_img_width; j <<= 1, j2++) - { - img_width = (1 << j2); - - img_format.image_channel_order = CL_RGBA; - img_format.image_channel_data_type = CL_UNORM_INT8; - streams[0] = - create_image_2d(context, CL_MEM_READ_WRITE, &img_format, - img_width, img_height, 0, NULL, NULL); - if (!streams[0]) - { - log_error("create_image_2d failed. width = %d, height = %d\n", img_width, img_height); - delete[] input_ptr; - delete[] output_ptr; - free_mtdata(d); - return -1; - } - img_format.image_channel_order = CL_RGBA; - img_format.image_channel_data_type = CL_UNORM_INT8; - streams[1] = - create_image_2d(context, CL_MEM_READ_WRITE, &img_format, - img_width, img_height, 0, NULL, NULL); - if (!streams[1]) - { - log_error("create_image_2d failed. width = %d, height = %d\n", img_width, img_height); - clReleaseMemObject(streams[0]); - delete[] input_ptr; - delete[] output_ptr; - free_mtdata(d); - return -1; - } + clMemWrapper streams[2]; - size_t origin[3] = {0,0,0}; - size_t region[3] = {img_width, img_height, 1}; - err = clEnqueueWriteImage(queue, streams[0], CL_FALSE, origin, region, 0, 0, input_ptr, 0, NULL, NULL); - if (err != CL_SUCCESS) - { - log_error("clWriteImage failed\n"); - clReleaseMemObject(streams[0]); - clReleaseMemObject(streams[1]); - delete[] input_ptr; - delete[] output_ptr; - free_mtdata(d); - return -1; - } + std::vector input(4 * img_width * img_height); + std::vector output(4 * img_width * img_height); - err = clSetKernelArg(kernel, 0, sizeof streams[0], &streams[0]); - err |= clSetKernelArg(kernel, 1, sizeof streams[1], &streams[1]); - err |= clSetKernelArg(kernel, 2, sizeof sampler, &sampler); - if (err != CL_SUCCESS) - { - log_error("clSetKernelArgs failed\n"); - clReleaseMemObject(streams[0]); - clReleaseMemObject(streams[1]); - delete[] input_ptr; - delete[] output_ptr; - free_mtdata(d); - return -1; - } + generate_random_inputs(input); - threads[0] = (size_t)img_width; - threads[1] = (size_t)img_height; - log_info("Testing image dimensions %d x %d with local threads NULL.\n", img_width, img_height); - err = clEnqueueNDRangeKernel( queue, kernel, 2, NULL, threads, NULL, 0, NULL, NULL ); - if (err != CL_SUCCESS) - { - log_error("clEnqueueNDRangeKernel failed\n"); - log_error("Image Dimension test failed. image width = %d, image height = %d, local NULL\n", - img_width, img_height); - clReleaseMemObject(streams[0]); - clReleaseMemObject(streams[1]); - delete[] input_ptr; - delete[] output_ptr; - free_mtdata(d); - return -1; - } - err = clEnqueueReadImage(queue, streams[1], CL_TRUE, origin, region, 0, 0, output_ptr, 0, NULL, NULL); - if (err != CL_SUCCESS) - { - log_error("clReadImage failed\n"); - log_error("Image Dimension test failed. image width = %d, image height = %d, local NULL\n", - img_width, img_height); - clReleaseMemObject(streams[0]); - clReleaseMemObject(streams[1]); - delete[] input_ptr; - delete[] output_ptr; - free_mtdata(d); - return -1; - } - err = verify_8888_image(input_ptr, output_ptr, img_width, img_height); - if (err) - { - total_errors++; - log_error("Image Dimension test failed. image width = %d, image height = %d\n", img_width, img_height); - } + const cl_image_format img_format = { CL_RGBA, CL_UNORM_INT8 }; - clReleaseMemObject(streams[0]); - clReleaseMemObject(streams[1]); - } - } + streams[0] = create_image_2d(context, CL_MEM_READ_WRITE, &img_format, + img_width, img_height, 0, nullptr, &err); + test_error(err, "create_image_2d failed"); + + streams[1] = create_image_2d(context, CL_MEM_READ_WRITE, &img_format, + img_width, img_height, 0, nullptr, &err); + test_error(err, "create_image_2d failed"); + + size_t origin[3] = { 0, 0, 0 }; + size_t region[3] = { img_width, img_height, 1 }; + err = clEnqueueWriteImage(queue, streams[0], CL_FALSE, origin, region, 0, 0, + input.data(), 0, nullptr, nullptr); + test_error(err, "clEnqueueWriteImage failed"); - // cleanup - delete[] input_ptr; - delete[] output_ptr; - free_mtdata(d); - clReleaseSampler(sampler); - clReleaseKernel(kernel); - clReleaseProgram(program); + clSamplerWrapper sampler = clCreateSampler( + context, CL_FALSE, CL_ADDRESS_CLAMP_TO_EDGE, CL_FILTER_NEAREST, &err); + test_error(err, "clCreateSampler failed"); + err = clSetKernelArg(kernel, 0, sizeof streams[0], &streams[0]); + err |= clSetKernelArg(kernel, 1, sizeof streams[1], &streams[1]); + err |= clSetKernelArg(kernel, 2, sizeof sampler, &sampler); + test_error(err, "clSetKernelArg failed"); + + size_t threads[] = { img_width, img_height }; + if (local_threads) + log_info( + "Testing image dimensions %d x %d with local threads %d x %d.\n", + img_width, img_height, local_threads[0], local_threads[1]); + else + log_info( + "Testing image dimensions %d x %d with local threads nullptr.\n", + img_width, img_height); + err = clEnqueueNDRangeKernel(queue, kernel, 2, nullptr, threads, + local_threads, 0, nullptr, nullptr); + test_error(err, "clEnqueueNDRangeKernel failed"); + + err = clEnqueueReadImage(queue, streams[1], CL_TRUE, origin, region, 0, 0, + output.data(), 0, nullptr, nullptr); + test_error(err, "clEnqueueReadImage failed"); + + if (0 != memcmp(input.data(), output.data(), 4 * img_width * img_height)) + { + total_errors++; + log_error("Image Dimension test failed. image width = %d, " + "image height = %d\n", + img_width, img_height); + } return total_errors; } +} +int test_imagedim_pow2(cl_device_id device, cl_context context, + cl_command_queue queue, int n_elems) +{ + clProgramWrapper program; + clKernelWrapper kernel; + size_t max_img_width; + size_t max_img_height; -int -test_imagedim_non_pow2(cl_device_id device, cl_context context, cl_command_queue queue, int n_elems) -{ - cl_mem streams[2]; - cl_image_format img_format; - unsigned char *input_ptr, *output_ptr; - cl_program program; - cl_kernel kernel; - size_t threads[2], local_threads[2]; - cl_ulong max_mem_size; - size_t img_width, max_img_width; - size_t img_height, max_img_height; - size_t max_img_dim; - int i, j, i2, j2, err = 0; - size_t max_image2d_width, max_image2d_height; + int err = 0; int total_errors = 0; - size_t max_local_workgroup_size[3]; - MTdata d; - - PASSIVE_REQUIRE_IMAGE_SUPPORT( device ) - err = create_single_kernel_helper( context, &program, &kernel, 1, &image_dim_kernel_code, "test_image_dim" ); - if (err) - { - log_error("create_program_and_kernel_with_sources failed\n"); - return -1; - } + PASSIVE_REQUIRE_IMAGE_SUPPORT(device) - size_t work_group_size = 0; - err = clGetKernelWorkGroupInfo(kernel, device, CL_KERNEL_WORK_GROUP_SIZE, sizeof(work_group_size), &work_group_size, NULL); - test_error(err, "clGetKerenlWorkgroupInfo failed for CL_KERNEL_WORK_GROUP_SIZE"); + err = create_single_kernel_helper(context, &program, &kernel, 1, + &image_dim_kernel_code, "test_image_dim"); + test_error(err, "create_single_kernel_helper failed"); - err = clGetDeviceInfo(device, CL_DEVICE_MAX_WORK_ITEM_SIZES, sizeof(max_local_workgroup_size), max_local_workgroup_size, NULL); - test_error(err, "clGetDeviceInfo failed for CL_DEVICE_MAX_WORK_ITEM_SIZES"); + err = get_max_image_dimensions(device, max_img_width, max_img_height); + test_error(err, "get_max_image_dimensions failed"); - err = clGetDeviceInfo(device, CL_DEVICE_GLOBAL_MEM_SIZE,sizeof(max_mem_size), &max_mem_size, NULL); - if (err) - { - log_error("clGetDeviceInfo for CL_DEVICE_GLOBAL_MEM_SIZE failed (%d)\n", err); - return -1; - } - err = clGetDeviceInfo(device, CL_DEVICE_IMAGE2D_MAX_WIDTH, sizeof(max_image2d_width), &max_image2d_width, NULL); - if (err) - { - log_error("clGetDeviceInfo for CL_DEVICE_IMAGE2D_MAX_WIDTH failed (%d)\n", err); - return -1; - } - err = clGetDeviceInfo(device, CL_DEVICE_IMAGE2D_MAX_HEIGHT, sizeof(max_image2d_width), &max_image2d_height, NULL); - if (err) + // test power of 2 width, height starting at 1 to 4K + for (size_t i = 1, i2 = 0; i <= max_img_height; i <<= 1, i2++) { - log_error("clGetDeviceInfo for CL_DEVICE_IMAGE2D_MAX_HEIGHT failed (%d)\n", err); - return -1; + size_t img_height = (1 << i2); + for (size_t j = 1, j2 = 0; j <= max_img_width; j <<= 1, j2++) + { + size_t img_width = (1 << j2); + + total_errors += test_imagedim_common( + context, queue, kernel, nullptr, img_width, img_height); + } } - log_info("Device reported max image sizes of %lu x %lu, and max mem size of %gMB.\n", - max_image2d_width, max_image2d_height, max_mem_size/(1024.0*1024.0)); - cl_sampler sampler = clCreateSampler(context, CL_FALSE, CL_ADDRESS_CLAMP_TO_EDGE, CL_FILTER_NEAREST, &err); - test_error(err, "clCreateSampler failed"); + return total_errors; +} - max_img_width = (int)max_image2d_width; - max_img_height = (int)max_image2d_height; - if (max_mem_size > (cl_ulong)SIZE_MAX) { - max_mem_size = (cl_ulong)SIZE_MAX; - } +int test_imagedim_non_pow2(cl_device_id device, cl_context context, + cl_command_queue queue, int n_elems) +{ + clProgramWrapper program; + clKernelWrapper kernel; - // determine max image dim we can allocate - assume RGBA image, 4 bytes per pixel, - // and we want to consume 1/4 of global memory (this is the minimum required to be - // supported by the spec) - max_mem_size /= 4; // use 1/4 - max_mem_size /= 4; // 4 bytes per pixel - max_img_dim = (int)sqrt((double)max_mem_size); - // convert to a power of 2 - { - unsigned int n = (unsigned int)max_img_dim; - unsigned int m = 0x80000000; + size_t max_img_width; + size_t max_img_height; + size_t max_local_workgroup_size[3] = {}; + size_t work_group_size = 0; + int err = 0; + int total_errors = 0; - // round-down to the nearest power of 2 - while (m > n) - m >>= 1; - max_img_dim = (int)m; - } + PASSIVE_REQUIRE_IMAGE_SUPPORT(device) - if (max_img_width > max_img_dim) - max_img_width = max_img_dim; - if (max_img_height > max_img_dim) - max_img_height = max_img_dim; + err = create_single_kernel_helper(context, &program, &kernel, 1, + &image_dim_kernel_code, "test_image_dim"); + test_error(err, "create_single_kernel_helper failed"); - log_info("Adjusted maximum image size to test is %d x %d, which is a max mem size of %gMB.\n", - max_img_width, max_img_height, (max_img_width*max_img_height*4)/(1024.0*1024.0)); + err = get_max_image_dimensions(device, max_img_width, max_img_height); + test_error(err, "get_max_image_dimensions failed"); + + err = clGetKernelWorkGroupInfo(kernel, device, CL_KERNEL_WORK_GROUP_SIZE, + sizeof(work_group_size), &work_group_size, + nullptr); + test_error(err, + "clGetKernelWorkgroupInfo failed for CL_KERNEL_WORK_GROUP_SIZE"); + + err = clGetDeviceInfo(device, CL_DEVICE_MAX_WORK_ITEM_SIZES, + sizeof(max_local_workgroup_size), + max_local_workgroup_size, nullptr); + test_error(err, "clGetDeviceInfo failed for CL_DEVICE_MAX_WORK_ITEM_SIZES"); - d = init_genrand( gRandomSeed ); - input_ptr = generate_8888_image(max_img_width, max_img_height, d); - output_ptr = new unsigned char[4 * max_img_width * max_img_height]; + // clamp max_local_workgroup_size to CL_KERNEL_WORK_GROUP_SIZE + for (auto &max_lws : max_local_workgroup_size) + max_lws = std::min(max_lws, work_group_size); - int plus_minus; - for (plus_minus = 0; plus_minus < 3; plus_minus++) + for (int plus_minus = 0; plus_minus < 3; plus_minus++) { - // test power of 2 width, height starting at 1 to 4K - for (i=2,i2=1; i<=max_img_height; i<<=1,i2++) + // test power of 2 width, height starting at 1 to 4K + for (size_t i = 2, i2 = 1; i <= max_img_height; i <<= 1, i2++) { - img_height = (1 << i2); - for (j=2,j2=1; j<=max_img_width; j<<=1,j2++) + size_t img_height = (1 << i2); + for (size_t j = 2, j2 = 1; j <= max_img_width; j <<= 1, j2++) { - img_width = (1 << j2); + size_t img_width = (1 << j2); size_t effective_img_height = img_height; size_t effective_img_width = img_width; - local_threads[0] = 1; - local_threads[1] = 1; + size_t local_threads[] = { 1, 1 }; - switch (plus_minus) { + switch (plus_minus) + { case 0: - effective_img_height--; - local_threads[0] = work_group_size > max_local_workgroup_size[0] ? max_local_workgroup_size[0] : work_group_size; - while (img_width%local_threads[0] != 0) - local_threads[0]--; - break; + effective_img_height--; + local_threads[0] = max_local_workgroup_size[0]; + while (img_width % local_threads[0] != 0) + local_threads[0]--; + break; case 1: - effective_img_width--; - local_threads[1] = work_group_size > max_local_workgroup_size[1] ? max_local_workgroup_size[1] : work_group_size; - while (img_height%local_threads[1] != 0) - local_threads[1]--; - break; + effective_img_width--; + local_threads[1] = max_local_workgroup_size[1]; + while (img_height % local_threads[1] != 0) + local_threads[1]--; + break; case 2: - effective_img_width--; - effective_img_height--; - break; - default: - break; - } - - img_format.image_channel_order = CL_RGBA; - img_format.image_channel_data_type = CL_UNORM_INT8; - streams[0] = create_image_2d( - context, CL_MEM_READ_WRITE, &img_format, - effective_img_width, effective_img_height, 0, NULL, NULL); - if (!streams[0]) - { - log_error("create_image_2d failed. width = %d, height = %d\n", effective_img_width, effective_img_height); - delete[] input_ptr; - delete[] output_ptr; - free_mtdata(d); - return -1; - } - img_format.image_channel_order = CL_RGBA; - img_format.image_channel_data_type = CL_UNORM_INT8; - streams[1] = create_image_2d( - context, CL_MEM_READ_WRITE, &img_format, - effective_img_width, effective_img_height, 0, NULL, NULL); - if (!streams[1]) - { - log_error("create_image_2d failed. width = %d, height = %d\n", effective_img_width, effective_img_height); - clReleaseMemObject(streams[0]); - delete[] input_ptr; - delete[] output_ptr; - free_mtdata(d); - return -1; - } - - size_t origin[3] = {0,0,0}; - size_t region[3] = {effective_img_width, effective_img_height, 1}; - err = clEnqueueWriteImage(queue, streams[0], CL_FALSE, origin, region, 0, 0, input_ptr, 0, NULL, NULL); - if (err != CL_SUCCESS) - { - log_error("clWriteImage failed\n"); - clReleaseMemObject(streams[0]); - clReleaseMemObject(streams[1]); - delete[] input_ptr; - delete[] output_ptr; - free_mtdata(d); - return -1; - } - - err = clSetKernelArg(kernel, 0, sizeof streams[0], &streams[0]); - err |= clSetKernelArg(kernel, 1, sizeof streams[1], &streams[1]); - err |= clSetKernelArg(kernel, 2, sizeof sampler, &sampler); - if (err != CL_SUCCESS) - { - log_error("clSetKernelArgs failed\n"); - clReleaseMemObject(streams[0]); - clReleaseMemObject(streams[1]); - delete[] input_ptr; - delete[] output_ptr; - free_mtdata(d); - return -1; - } - - threads[0] = (size_t)effective_img_width; - threads[1] = (size_t)effective_img_height; - log_info("Testing image dimensions %d x %d with local threads %d x %d.\n", - effective_img_width, effective_img_height, (int)local_threads[0], (int)local_threads[1]); - err = clEnqueueNDRangeKernel( queue, kernel, 2, NULL, threads, local_threads, 0, NULL, NULL ); - if (err != CL_SUCCESS) - { - log_error("clEnqueueNDRangeKernel failed\n"); - log_error("Image Dimension test failed. image width = %d, image height = %d, local %d x %d\n", - effective_img_width, effective_img_height, (int)local_threads[0], (int)local_threads[1]); - clReleaseMemObject(streams[0]); - clReleaseMemObject(streams[1]); - delete[] input_ptr; - delete[] output_ptr; - free_mtdata(d); - return -1; - } - err = clEnqueueReadImage(queue, streams[1], CL_TRUE, origin, region, 0, 0, output_ptr, 0, NULL, NULL); - if (err != CL_SUCCESS) - { - log_error("clReadImage failed\n"); - log_error("Image Dimension test failed. image width = %d, image height = %d, local %d x %d\n", - effective_img_width, effective_img_height, (int)local_threads[0], (int)local_threads[1]); - clReleaseMemObject(streams[0]); - clReleaseMemObject(streams[1]); - delete[] input_ptr; - delete[] output_ptr; - free_mtdata(d); - return -1; - } - err = verify_8888_image(input_ptr, output_ptr, effective_img_width, effective_img_height); - if (err) - { - total_errors++; - log_error("Image Dimension test failed. image width = %d, image height = %d\n", effective_img_width, effective_img_height); + effective_img_width--; + effective_img_height--; + break; + default: break; } - clReleaseMemObject(streams[0]); - clReleaseMemObject(streams[1]); + total_errors += test_imagedim_common( + context, queue, kernel, local_threads, effective_img_width, + effective_img_height); } } + } - } - - // cleanup - delete[] input_ptr; - delete[] output_ptr; - free_mtdata(d); - clReleaseSampler(sampler); - clReleaseKernel(kernel); - clReleaseProgram(program); - - return total_errors; + return total_errors; } - - - - diff --git a/test_conformance/computeinfo/main.cpp b/test_conformance/computeinfo/main.cpp index e382b38e4c..b6350c1c2a 100644 --- a/test_conformance/computeinfo/main.cpp +++ b/test_conformance/computeinfo/main.cpp @@ -1452,5 +1452,9 @@ int main(int argc, const char** argv) } } - return runTestHarness(argCount, argList, test_num, test_list, true, 0); + int error = runTestHarness(argCount, argList, test_num, test_list, true, 0); + + free(argList); + + return error; } diff --git a/test_conformance/extensions/cl_khr_external_semaphore/CMakeLists.txt b/test_conformance/extensions/cl_khr_external_semaphore/CMakeLists.txt index 6e02ba97e5..df136004cd 100644 --- a/test_conformance/extensions/cl_khr_external_semaphore/CMakeLists.txt +++ b/test_conformance/extensions/cl_khr_external_semaphore/CMakeLists.txt @@ -15,12 +15,7 @@ include_directories(${CMAKE_CURRENT_SOURCE_DIR}) include_directories (${CLConform_INCLUDE_DIR}) -# needed by Vulkan wrapper to link -if(WIN32) - list(APPEND CLConform_LIBRARIES vulkan-1 vulkan_wrapper) -else(WIN32) - list(APPEND CLConform_LIBRARIES vulkan dl vulkan_wrapper) -endif(WIN32) +list(APPEND CLConform_LIBRARIES vulkan_wrapper) set(CMAKE_CXX_FLAGS "-fpermissive") include_directories("../../common/vulkan_wrapper") diff --git a/test_conformance/extensions/cl_khr_semaphore/CMakeLists.txt b/test_conformance/extensions/cl_khr_semaphore/CMakeLists.txt index ed359d5983..b3bdd34609 100644 --- a/test_conformance/extensions/cl_khr_semaphore/CMakeLists.txt +++ b/test_conformance/extensions/cl_khr_semaphore/CMakeLists.txt @@ -3,6 +3,7 @@ set(MODULE_NAME CL_KHR_SEMAPHORE) set(${MODULE_NAME}_SOURCES main.cpp test_semaphores.cpp + test_semaphores_negative_getinfo.cpp test_semaphores_negative_wait.cpp test_semaphores_negative_create.cpp semaphore_base.h diff --git a/test_conformance/extensions/cl_khr_semaphore/main.cpp b/test_conformance/extensions/cl_khr_semaphore/main.cpp index de996a9484..ec482ce309 100644 --- a/test_conformance/extensions/cl_khr_semaphore/main.cpp +++ b/test_conformance/extensions/cl_khr_semaphore/main.cpp @@ -46,6 +46,9 @@ test_definition test_list[] = { ADD_TEST_VERSION(semaphores_negative_create_invalid_value, Version(1, 2)), ADD_TEST_VERSION(semaphores_negative_create_invalid_operation, Version(1, 2)), + ADD_TEST_VERSION(semaphores_negative_get_info_invalid_semaphore, + Version(1, 2)), + ADD_TEST_VERSION(semaphores_negative_get_info_invalid_value, Version(1, 2)), ADD_TEST_VERSION(semaphores_negative_wait_invalid_command_queue, Version(1, 2)), ADD_TEST_VERSION(semaphores_negative_wait_invalid_value, Version(1, 2)), diff --git a/test_conformance/extensions/cl_khr_semaphore/procs.h b/test_conformance/extensions/cl_khr_semaphore/procs.h index aeb601b6e6..87a95d38d7 100644 --- a/test_conformance/extensions/cl_khr_semaphore/procs.h +++ b/test_conformance/extensions/cl_khr_semaphore/procs.h @@ -67,6 +67,12 @@ extern int test_semaphores_negative_create_invalid_value(cl_device_id device, extern int test_semaphores_negative_create_invalid_operation( cl_device_id device, cl_context context, cl_command_queue queue, int num_elements); +extern int test_semaphores_negative_get_info_invalid_semaphore( + cl_device_id device, cl_context context, cl_command_queue queue, + int num_elements); +extern int test_semaphores_negative_get_info_invalid_value( + cl_device_id device, cl_context context, cl_command_queue queue, + int num_elements); extern int test_semaphores_negative_wait_invalid_command_queue( cl_device_id device, cl_context context, cl_command_queue queue, int num_elements); diff --git a/test_conformance/extensions/cl_khr_semaphore/test_semaphores_negative_getinfo.cpp b/test_conformance/extensions/cl_khr_semaphore/test_semaphores_negative_getinfo.cpp new file mode 100644 index 0000000000..0cf8bb0faf --- /dev/null +++ b/test_conformance/extensions/cl_khr_semaphore/test_semaphores_negative_getinfo.cpp @@ -0,0 +1,130 @@ +// +// Copyright (c) 2024 The Khronos Group Inc. +// +// Licensed under the Apache License, Version 2.0 (the "License"); +// you may not use this file except in compliance with the License. +// You may obtain a copy of the License at +// +// http://www.apache.org/licenses/LICENSE-2.0 +// +// Unless required by applicable law or agreed to in writing, software +// distributed under the License is distributed on an "AS IS" BASIS, +// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +// See the License for the specific language governing permissions and +// limitations under the License. +// + +#include "semaphore_base.h" + +#include "harness/errorHelpers.h" + +namespace { + +// sema_object is not a valid semaphore. + +struct GetInfoInvalidSemaphore : public SemaphoreTestBase +{ + GetInfoInvalidSemaphore(cl_device_id device, cl_context context, + cl_command_queue queue) + : SemaphoreTestBase(device, context, queue) + {} + + cl_int Run() override + { + // Wait semaphore + cl_semaphore_type_khr type = 0; + size_t ret_size = 0; + cl_int err = clGetSemaphoreInfoKHR(nullptr, CL_SEMAPHORE_TYPE_KHR, + sizeof(cl_semaphore_type_khr), &type, + &ret_size); + test_failure_error(err, CL_INVALID_SEMAPHORE_KHR, + "Unexpected clGetSemaphoreInfoKHR return"); + + return CL_SUCCESS; + } +}; + +// 1) param_name is not one of the attribute defined in the Semaphore Queries +// table + +// 2) param_value_size is less than the size of Return Type of the corresponding +// param_name attribute as defined in the Semaphore Queries table. + +struct GetInfoInvalidValue : public SemaphoreTestBase +{ + GetInfoInvalidValue(cl_device_id device, cl_context context, + cl_command_queue queue) + : SemaphoreTestBase(device, context, queue) + {} + + cl_int Run() override + { + // Create semaphore + cl_semaphore_properties_khr sema_props[] = { + static_cast(CL_SEMAPHORE_TYPE_KHR), + static_cast( + CL_SEMAPHORE_TYPE_BINARY_KHR), + static_cast( + CL_SEMAPHORE_DEVICE_HANDLE_LIST_KHR), + (cl_semaphore_properties_khr)device, + CL_SEMAPHORE_DEVICE_HANDLE_LIST_END_KHR, + 0 + }; + + cl_int err = CL_SUCCESS; + semaphore = + clCreateSemaphoreWithPropertiesKHR(context, sema_props, &err); + test_error(err, "Could not create semaphore"); + + // (1) + cl_semaphore_info_khr param_name = ~0; + err = clGetSemaphoreInfoKHR(semaphore, param_name, 0, nullptr, nullptr); + test_failure_error(err, CL_INVALID_VALUE, + "Unexpected clGetSemaphoreInfoKHR return"); + + // (2) + size_t size = 0; + err = clGetSemaphoreInfoKHR(semaphore, CL_SEMAPHORE_PROPERTIES_KHR, 0, + nullptr, &size); + test_error(err, "Could not query semaphore"); + + // make sure that first test provides too small param size + if (size != sizeof(sema_props)) + test_fail("Error: expected size %d, returned %d", + sizeof(sema_props), size); + + // first test with non-zero property size but not enough + cl_semaphore_properties_khr ret_props = 0; + err = clGetSemaphoreInfoKHR(semaphore, CL_SEMAPHORE_PROPERTIES_KHR, + sizeof(ret_props), &ret_props, nullptr); + test_failure_error(err, CL_INVALID_VALUE, + "Unexpected clGetSemaphoreInfoKHR return"); + + // second test with zero property size + cl_semaphore_type_khr type = 0; + err = clGetSemaphoreInfoKHR(semaphore, CL_SEMAPHORE_TYPE_KHR, 0, &type, + nullptr); + test_failure_error(err, CL_INVALID_VALUE, + "Unexpected clGetSemaphoreInfoKHR return"); + + return CL_SUCCESS; + } +}; + +} + +int test_semaphores_negative_get_info_invalid_semaphore(cl_device_id device, + cl_context context, + cl_command_queue queue, + int num_elements) +{ + return MakeAndRunTest(device, context, queue); +} + +int test_semaphores_negative_get_info_invalid_value(cl_device_id device, + cl_context context, + cl_command_queue queue, + int num_elements) +{ + return MakeAndRunTest(device, context, queue); +} diff --git a/test_conformance/math_brute_force/function_list.cpp b/test_conformance/math_brute_force/function_list.cpp index 832615e1bc..0b4ec0e544 100644 --- a/test_conformance/math_brute_force/function_list.cpp +++ b/test_conformance/math_brute_force/function_list.cpp @@ -271,23 +271,19 @@ const Func functionList[] = { ENTRY(cosh, 4.0f, 4.0f, 2.f, FTZ_OFF, unaryF), ENTRY_EXT(cospi, 4.0f, 4.0f, 2.f, 0.00048828125f, FTZ_OFF, unaryF, 0.00048828125f), // relaxed ulp 2^-11 - // ENTRY( erfc, 16.0f, - // 16.0f, FTZ_OFF, unaryF), - // //disabled for 1.0 due to lack of - // reference implementation ENTRY( erf, - // 16.0f, 16.0f, FTZ_OFF, - // unaryF), //disabled for 1.0 due to lack - // of reference implementation - ENTRY_EXT(exp, 3.0f, 4.0f, 2.f, 3.0f, FTZ_OFF, unaryF, - 4.0f), // relaxed error is actually overwritten in unary.c as it - // is 3+floor(fabs(2*x)) - ENTRY_EXT(exp2, 3.0f, 4.0f, 2.f, 3.0f, FTZ_OFF, unaryF, - 4.0f), // relaxed error is actually overwritten in unary.c as it - // is 3+floor(fabs(2*x)) - ENTRY_EXT(exp10, 3.0f, 4.0f, 2.f, 8192.0f, FTZ_OFF, unaryF, - 8192.0f), // relaxed error is actually overwritten in unary.c as - // it is 3+floor(fabs(2*x)) in derived mode, + //ENTRY(erfc, 16.0f, 16.0f, FTZ_OFF, unaryF), //disabled for 1.0 due to lack of reference implementation + //ENTRY(erf, 16.0f, 16.0f, FTZ_OFF, unaryF), //disabled for 1.0 due to lack of reference implementation + + // relaxed error is overwritten in unary.c as it is 3+floor(fabs(2*x)) + ENTRY_EXT(exp, 3.0f, 4.0f, 2.f, 3.0f, FTZ_OFF, unaryF, 4.0f), + + // relaxed error is overwritten in unary.c as it is 3+floor(fabs(2*x)) + ENTRY_EXT(exp2, 3.0f, 4.0f, 2.f, 3.0f, FTZ_OFF, unaryF, 4.0f), + + // relaxed error is overwritten in unary.c as it is 3+floor(fabs(2*x)) in derived mode; // in non-derived mode it uses the ulp error for half_exp10. + ENTRY_EXT(exp10, 3.0f, 4.0f, 2.f, 8192.0f, FTZ_OFF, unaryF, 8192.0f), + ENTRY(expm1, 3.0f, 4.0f, 2.f, FTZ_OFF, unaryF), ENTRY(fabs, 0.0f, 0.0f, 0.0f, FTZ_OFF, unaryF), ENTRY(fdim, 0.0f, 0.0f, 0.0f, FTZ_OFF, binaryF), @@ -325,23 +321,23 @@ const Func functionList[] = { 4.76837158203125e-7f), // relaxed ulp 2^-21 ENTRY(log1p, 2.0f, 4.0f, 2.0f, FTZ_OFF, unaryF), ENTRY(logb, 0.0f, 0.0f, 0.0f, FTZ_OFF, unaryF), - ENTRY_EXT(mad, INFINITY, INFINITY, INFINITY, INFINITY, FTZ_OFF, - mad_function, - INFINITY), // in fast-relaxed-math mode it has to be either - // exactly rounded fma or exactly rounded a*b+c + + // In fast-relaxed-math mode it has to be either exactly rounded fma or exactly rounded a*b+c + ENTRY_EXT(mad, INFINITY, INFINITY, INFINITY, INFINITY, FTZ_OFF, mad_function, INFINITY), + ENTRY(maxmag, 0.0f, 0.0f, 0.0f, FTZ_OFF, binaryF), ENTRY(minmag, 0.0f, 0.0f, 0.0f, FTZ_OFF, binaryF), ENTRY(modf, 0.0f, 0.0f, 0.0f, FTZ_OFF, unaryF_two_results), ENTRY(nan, 0.0f, 0.0f, 0.0f, FTZ_OFF, unaryF_u), ENTRY(nextafter, 0.0f, 0.0f, 0.0f, FTZ_OFF, binaryF_nextafter), - ENTRY_EXT(pow, 16.0f, 16.0f, 4.0f, 8192.0f, FTZ_OFF, binaryF, - 8192.0f), // in derived mode the ulp error is calculated as - // exp2(y*log2(x)) and in non-derived it is the same as - // half_pow + + // In derived mode the ulp error is calculated as exp2(y*log2(x)). + // In non-derived it is the same as half_pow. + ENTRY_EXT(pow, 16.0f, 16.0f, 4.0f, 8192.0f, FTZ_OFF, binaryF, 8192.0f), + ENTRY(pown, 16.0f, 16.0f, 4.0f, FTZ_OFF, binaryF_i), ENTRY(powr, 16.0f, 16.0f, 4.0f, FTZ_OFF, binaryF), - // ENTRY( reciprocal, 1.0f, - // 1.0f, FTZ_OFF, unaryF), + //ENTRY(reciprocal, 1.0f, 1.0f, FTZ_OFF, unaryF), ENTRY(remainder, 0.0f, 0.0f, 0.0f, FTZ_OFF, binaryF), ENTRY(remquo, 0.0f, 0.0f, 0.0f, FTZ_OFF, binaryF_two_results_i), ENTRY(rint, 0.0f, 0.0f, 0.0f, FTZ_OFF, unaryF), @@ -385,16 +381,14 @@ const Func functionList[] = { FTZ_OFF, RELAXED_OFF, unaryF }, - ENTRY_EXT( - tan, 5.0f, 5.0f, 2.0f, 8192.0f, FTZ_OFF, unaryF, - 8192.0f), // in derived mode it the ulp error is calculated as sin/cos - // and in non-derived mode it is the same as half_tan. + + // In derived mode it the ulp error is calculated as sin/cos. + // In non-derived mode it is the same as half_tan. + ENTRY_EXT(tan, 5.0f, 5.0f, 2.0f, 8192.0f, FTZ_OFF, unaryF, 8192.0f), + ENTRY(tanh, 5.0f, 5.0f, 2.0f, FTZ_OFF, unaryF), ENTRY(tanpi, 6.0f, 6.0f, 2.0f, FTZ_OFF, unaryF), - // ENTRY( tgamma, 16.0f, - // 16.0f, FTZ_OFF, unaryF), - // // Commented this out until we can be - // sure this requirement is realistic + //ENTRY(tgamma, 16.0f, 16.0f, FTZ_OFF, unaryF), Commented this out until we can be sure this requirement is realistic ENTRY(trunc, 0.0f, 0.0f, 0.0f, FTZ_OFF, unaryF), HALF_ENTRY(cos, 8192.0f, 8192.0f, FTZ_ON, unaryOF), diff --git a/test_conformance/printf/test_printf.cpp b/test_conformance/printf/test_printf.cpp index 2ecf400180..3d539ed572 100644 --- a/test_conformance/printf/test_printf.cpp +++ b/test_conformance/printf/test_printf.cpp @@ -15,11 +15,17 @@ // #include "harness/os_helpers.h" #include "harness/typeWrappers.h" +#include "harness/stringHelpers.h" +#include "harness/conversions.h" +#include +#include #include -#include +#include #include #include +#include +#include #if ! defined( _WIN32) #if defined(__APPLE__) @@ -43,6 +49,7 @@ #include "harness/errorHelpers.h" #include "harness/kernelHelpers.h" #include "harness/parseParameters.h" +#include "harness/rounding_mode.h" #include @@ -51,50 +58,49 @@ typedef unsigned int uint32_t; test_status InitCL( cl_device_id device ); +namespace { + //----------------------------------------- -// Static helper functions declaration +// helper functions declaration //----------------------------------------- -static void printUsage( void ); - //Stream helper functions //Associate stdout stream with the file(gFileName):i.e redirect stdout stream to the specific files (gFileName) -static int acquireOutputStream(int* error); +int acquireOutputStream(int* error); //Close the file(gFileName) associated with the stdout stream and disassociates it. -static void releaseOutputStream(int fd); +void releaseOutputStream(int fd); //Get analysis buffer to verify the correctess of printed data -static void getAnalysisBuffer(char* analysisBuffer); +void getAnalysisBuffer(char* analysisBuffer); //Kernel builder helper functions //Check if the test case is for kernel that has argument -static int isKernelArgument(testCase* pTestCase,size_t testId); +int isKernelArgument(testCase* pTestCase, size_t testId); //Check if the test case treats %p format for void* -static int isKernelPFormat(testCase* pTestCase,size_t testId); +int isKernelPFormat(testCase* pTestCase, size_t testId); //----------------------------------------- // Static functions declarations //----------------------------------------- // Make a program that uses printf for the given type/format, -static cl_program -makePrintfProgram(cl_kernel* kernel_ptr, const cl_context context, - const unsigned int testId, const unsigned int testNum, - const unsigned int formatNum, bool isLongSupport = true, - bool is64bAddrSpace = false); +cl_program makePrintfProgram(cl_kernel* kernel_ptr, const cl_context context, + cl_device_id device, const unsigned int testId, + const unsigned int testNum, + const unsigned int formatNum); // Creates and execute the printf test for the given device, context, type/format -static int doTest(cl_command_queue queue, cl_context context, - const unsigned int testId, cl_device_id device); +int doTest(cl_command_queue queue, cl_context context, + const unsigned int testId, cl_device_id device); // Check if device supports long -static bool isLongSupported(cl_device_id device_id); +bool isLongSupported(cl_device_id device_id); // Check if device address space is 64 bits -static bool is64bAddressSpace(cl_device_id device_id); +bool is64bAddressSpace(cl_device_id device_id); //Wait until event status is CL_COMPLETE int waitForEvent(cl_event* event); @@ -111,21 +117,25 @@ int s_test_cnt = 0; int s_test_fail = 0; int s_test_skip = 0; +cl_context gContext; +cl_command_queue gQueue; +int gFd; + +char gFileName[256]; -static cl_context gContext; -static cl_command_queue gQueue; -static int gFd; +MTdataHolder gMTdata; -static char gFileName[256]; +// For the sake of proper logging of negative results +std::string gLatestKernelSource; //----------------------------------------- -// Static helper functions definition +// helper functions definition //----------------------------------------- //----------------------------------------- // acquireOutputStream //----------------------------------------- -static int acquireOutputStream(int* error) +int acquireOutputStream(int* error) { int fd = streamDup(fileno(stdout)); *error = 0; @@ -140,7 +150,7 @@ static int acquireOutputStream(int* error) //----------------------------------------- // releaseOutputStream //----------------------------------------- -static void releaseOutputStream(int fd) +void releaseOutputStream(int fd) { fflush(stdout); streamDup2(fd,fileno(stdout)); @@ -150,7 +160,8 @@ static void releaseOutputStream(int fd) //----------------------------------------- // printfCallBack //----------------------------------------- -static void CL_CALLBACK printfCallBack(const char *printf_data, size_t len, size_t final, void *user_data) +void CL_CALLBACK printfCallBack(const char* printf_data, size_t len, + size_t final, void* user_data) { fwrite(printf_data, 1, len, stdout); } @@ -158,7 +169,7 @@ static void CL_CALLBACK printfCallBack(const char *printf_data, size_t len, size //----------------------------------------- // getAnalysisBuffer //----------------------------------------- -static void getAnalysisBuffer(char* analysisBuffer) +void getAnalysisBuffer(char* analysisBuffer) { FILE *fp; memset(analysisBuffer,0,ANALYSIS_BUFFER_SIZE); @@ -177,14 +188,14 @@ static void getAnalysisBuffer(char* analysisBuffer) //----------------------------------------- // isKernelArgument //----------------------------------------- -static int isKernelArgument(testCase* pTestCase,size_t testId) +int isKernelArgument(testCase* pTestCase, size_t testId) { return strcmp(pTestCase->_genParameters[testId].addrSpaceArgumentTypeQualifier,""); } //----------------------------------------- // isKernelPFormat //----------------------------------------- -static int isKernelPFormat(testCase* pTestCase,size_t testId) +int isKernelPFormat(testCase* pTestCase, size_t testId) { return strcmp(pTestCase->_genParameters[testId].addrSpacePAdd,""); } @@ -211,18 +222,159 @@ int waitForEvent(cl_event* event) } //----------------------------------------- -// Static helper functions definition +// makeMixedFormatPrintfProgram +// Generates in-flight printf kernel with format string including: +// -data before conversion flags (randomly generated ascii string) +// -randomly generated conversion flags (integer or floating point) +// -data after conversion flags (randomly generated ascii string). +// Moreover it generates suitable arguments. +// example: printf("zH, %u, %a, D+{gy\n", -929240879, 24295.671875f) //----------------------------------------- +cl_program makeMixedFormatPrintfProgram(cl_kernel* kernel_ptr, + const cl_context context, + const cl_device_id device, + const unsigned int testId, + const unsigned int testNum, + const std::string& testname) +{ + auto gen_char = [&]() { + static const char dict[] = { + " \t!#$&()*+,-./" + "123456789:;<=>?@ABCDEFGHIJKLMNOPQRSTUVWXYZ[]^_`" + "abcdefghijklmnopqrstuvwxyz{|}~" + }; + return dict[genrand_int32(gMTdata) % ((int)sizeof(dict) - 1)]; + }; + + std::array, 2> formats = { + { { "%f", "%e", "%g", "%a", "%F", "%E", "%G", "%A" }, + { "%d", "%i", "%u", "%x", "%o", "%X" } } + }; + std::vector data_before(2 + genrand_int32(gMTdata) % 8); + std::vector data_after(2 + genrand_int32(gMTdata) % 8); + + std::generate(data_before.begin(), data_before.end(), gen_char); + std::generate(data_after.begin(), data_after.end(), gen_char); + + cl_uint num_args = 2 + genrand_int32(gMTdata) % 4; + + // Map device rounding to CTS rounding type + // get_default_rounding_mode supports RNE and RTZ + auto get_rounding = [](const cl_device_fp_config& fpConfig) { + if (fpConfig == CL_FP_ROUND_TO_NEAREST) + { + return kRoundToNearestEven; + } + else if (fpConfig == CL_FP_ROUND_TO_ZERO) + { + return kRoundTowardZero; + } + else + { + assert(false && "Unreachable"); + } + return kDefaultRoundingMode; + }; + + const RoundingMode hostRound = get_round(); + RoundingMode deviceRound = get_rounding(get_default_rounding_mode(device)); + + std::ostringstream format_str; + std::ostringstream ref_str; + std::ostringstream source_gen; + std::ostringstream args_str; + source_gen << "__kernel void " << testname + << "(void)\n" + "{\n" + " printf(\""; + for (auto it : data_before) + { + format_str << it; + ref_str << it; + } + format_str << ", "; + ref_str << ", "; + + + for (cl_uint i = 0; i < num_args; i++) + { + std::uint8_t is_int = genrand_int32(gMTdata) % 2; + + // Set CPU rounding mode to match that of the device + set_round(deviceRound, is_int != 0 ? kint : kfloat); + + std::string format = + formats[is_int][genrand_int32(gMTdata) % formats[is_int].size()]; + format_str << format << ", "; + + if (is_int) + { + int arg = genrand_int32(gMTdata); + args_str << str_sprintf("%d", arg) << ", "; + ref_str << str_sprintf(format, arg) << ", "; + } + else + { + const float max_range = 100000.f; + float arg = get_random_float(-max_range, max_range, gMTdata); + args_str << str_sprintf("%f", arg) << "f, "; + ref_str << str_sprintf(format, arg) << ", "; + } + } + // Restore the original CPU rounding mode + set_round(hostRound, kfloat); + + for (auto it : data_after) + { + format_str << it; + ref_str << it; + } + + { + std::ostringstream args_cpy; + args_cpy << args_str.str(); + args_cpy.seekp(-2, std::ios_base::end); + args_cpy << ")\n"; + log_info("%d) testing printf(\"%s\\n\", %s", testNum, + format_str.str().c_str(), args_cpy.str().c_str()); + } + + args_str.seekp(-2, std::ios_base::end); + args_str << ");\n}\n"; + + + source_gen << format_str.str() << "\\n\"" + << ", " << args_str.str(); + + std::string kernel_source = source_gen.str(); + const char* ptr = kernel_source.c_str(); + + cl_program program; + cl_int err = create_single_kernel_helper(context, &program, kernel_ptr, 1, + &ptr, testname.c_str()); + + gLatestKernelSource = kernel_source.c_str(); + + // Save the reference result + allTestCase[testId]->_correctBuffer.push_back(ref_str.str()); + + if (!program || err) + { + log_error("create_single_kernel_helper failed\n"); + return NULL; + } + + return program; +} //----------------------------------------- // makePrintfProgram //----------------------------------------- -static cl_program makePrintfProgram(cl_kernel* kernel_ptr, - const cl_context context, - const unsigned int testId, - const unsigned int testNum, - const unsigned int formatNum, - bool isLongSupport, bool is64bAddrSpace) +cl_program makePrintfProgram(cl_kernel* kernel_ptr, const cl_context context, + const cl_device_id device, + const unsigned int testId, + const unsigned int testNum, + const unsigned int formatNum) { int err; cl_program program; @@ -293,6 +445,9 @@ static cl_program makePrintfProgram(cl_kernel* kernel_ptr, err = create_single_kernel_helper( context, &program, kernel_ptr, sizeof(sourceVec) / sizeof(sourceVec[0]), sourceVec, testname); + + gLatestKernelSource = + concat_kernel(sourceVec, sizeof(sourceVec) / sizeof(sourceVec[0])); } else if(allTestCase[testId]->_type == TYPE_ADDRESS_SPACE) { @@ -322,6 +477,15 @@ static cl_program makePrintfProgram(cl_kernel* kernel_ptr, sizeof(sourceAddrSpace) / sizeof(sourceAddrSpace[0]), sourceAddrSpace, testname); + + gLatestKernelSource = + concat_kernel(sourceAddrSpace, + sizeof(sourceAddrSpace) / sizeof(sourceAddrSpace[0])); + } + else if (allTestCase[testId]->_type == TYPE_MIXED_FORMAT_RANDOM) + { + return makeMixedFormatPrintfProgram(kernel_ptr, context, device, testId, + testNum, testname); } else { @@ -352,6 +516,8 @@ static cl_program makePrintfProgram(cl_kernel* kernel_ptr, err = create_single_kernel_helper(context, &program, kernel_ptr, 1, &ptr, testname); + + gLatestKernelSource = kernel_source.c_str(); } if (!program || err) { @@ -365,7 +531,7 @@ static cl_program makePrintfProgram(cl_kernel* kernel_ptr, //----------------------------------------- // isLongSupported //----------------------------------------- -static bool isLongSupported(cl_device_id device_id) +bool isLongSupported(cl_device_id device_id) { size_t tempSize = 0; cl_int status; @@ -409,7 +575,7 @@ static bool isLongSupported(cl_device_id device_id) //----------------------------------------- // is64bAddressSpace //----------------------------------------- -static bool is64bAddressSpace(cl_device_id device_id) +bool is64bAddressSpace(cl_device_id device_id) { cl_int status; cl_uint addrSpaceB; @@ -448,11 +614,78 @@ void subtest_fail(const char* msg, ...) ++s_test_cnt; } +//----------------------------------------- +// logTestType - printout test details +//----------------------------------------- + +void logTestType(const unsigned testId, const unsigned testNum, + unsigned formatNum) +{ + if (allTestCase[testId]->_type == TYPE_VECTOR) + { + log_info( + "%d)testing printf(\"%sv%s%s\",%s)\n", testNum, + allTestCase[testId]->_genParameters[testNum].vectorFormatFlag, + allTestCase[testId]->_genParameters[testNum].vectorSize, + allTestCase[testId]->_genParameters[testNum].vectorFormatSpecifier, + allTestCase[testId]->_genParameters[testNum].dataRepresentation); + } + else if (allTestCase[testId]->_type == TYPE_ADDRESS_SPACE) + { + if (isKernelArgument(allTestCase[testId], testNum)) + { + log_info("%d)testing kernel //argument %s \n printf(%s,%s)\n", + testNum, + allTestCase[testId] + ->_genParameters[testNum] + .addrSpaceArgumentTypeQualifier, + allTestCase[testId] + ->_genParameters[testNum] + .genericFormats[formatNum] + .c_str(), + allTestCase[testId] + ->_genParameters[testNum] + .addrSpaceParameter); + } + else + { + log_info("%d)testing kernel //variable %s \n printf(%s,%s)\n", + testNum, + allTestCase[testId] + ->_genParameters[testNum] + .addrSpaceVariableTypeQualifier, + allTestCase[testId] + ->_genParameters[testNum] + .genericFormats[formatNum] + .c_str(), + allTestCase[testId] + ->_genParameters[testNum] + .addrSpaceParameter); + } + } + else if (allTestCase[testId]->_type != TYPE_MIXED_FORMAT_RANDOM) + { + log_info("%d)testing printf(\"%s\"", testNum, + allTestCase[testId] + ->_genParameters[testNum] + .genericFormats[formatNum] + .c_str()); + if (allTestCase[testId]->_genParameters[testNum].dataRepresentation) + log_info(",%s", + allTestCase[testId] + ->_genParameters[testNum] + .dataRepresentation); + log_info(")\n"); + } + + fflush(stdout); +} + //----------------------------------------- // doTest //----------------------------------------- -static int doTest(cl_command_queue queue, cl_context context, - const unsigned int testId, cl_device_id device) +int doTest(cl_command_queue queue, cl_context context, + const unsigned int testId, cl_device_id device) { int err = TEST_FAIL; @@ -500,88 +733,13 @@ static int doTest(cl_command_queue queue, cl_context context, } } - for (unsigned formatNum = 0; formatNum < allTestCase[testId] - ->_genParameters[testNum] - .genericFormats.size(); + auto genParamsVec = allTestCase[testId]->_genParameters; + auto genFormatVec = genParamsVec[testNum].genericFormats; + + for (unsigned formatNum = 0; formatNum < genFormatVec.size(); formatNum++) { - if (allTestCase[testId]->_type == TYPE_VECTOR) - { - log_info( - "%d)testing printf(\"%sv%s%s\",%s)\n", testNum, - allTestCase[testId] - ->_genParameters[testNum] - .vectorFormatFlag, - allTestCase[testId]->_genParameters[testNum].vectorSize, - allTestCase[testId] - ->_genParameters[testNum] - .vectorFormatSpecifier, - allTestCase[testId] - ->_genParameters[testNum] - .dataRepresentation); - } - else if (allTestCase[testId]->_type == TYPE_ADDRESS_SPACE) - { - if (isKernelArgument(allTestCase[testId], testNum)) - { - log_info( - "%d)testing kernel //argument %s \n printf(%s,%s)\n", - testNum, - allTestCase[testId] - ->_genParameters[testNum] - .addrSpaceArgumentTypeQualifier, - allTestCase[testId] - ->_genParameters[testNum] - .genericFormats[formatNum] - .c_str(), - allTestCase[testId] - ->_genParameters[testNum] - .addrSpaceParameter); - } - else - { - log_info( - "%d)testing kernel //variable %s \n printf(%s,%s)\n", - testNum, - allTestCase[testId] - ->_genParameters[testNum] - .addrSpaceVariableTypeQualifier, - allTestCase[testId] - ->_genParameters[testNum] - .genericFormats[formatNum] - .c_str(), - allTestCase[testId] - ->_genParameters[testNum] - .addrSpaceParameter); - } - } - else - { - log_info("%d)testing printf(\"%s\"", testNum, - allTestCase[testId] - ->_genParameters[testNum] - .genericFormats[formatNum] - .c_str()); - if (allTestCase[testId] - ->_genParameters[testNum] - .dataRepresentation) - log_info(",%s", - allTestCase[testId] - ->_genParameters[testNum] - .dataRepresentation); - log_info(")\n"); - } - - fflush(stdout); - - // Long support for address in FULL_PROFILE/EMBEDDED_PROFILE - bool isLongSupport = true; - if (allTestCase[testId]->_type == TYPE_ADDRESS_SPACE - && isKernelPFormat(allTestCase[testId], testNum) - && !isLongSupported(device)) - { - isLongSupport = false; - } + logTestType(testId, testNum, formatNum); clProgramWrapper program; clKernelWrapper kernel; @@ -596,9 +754,8 @@ static int doTest(cl_command_queue queue, cl_context context, // execution. size_t globalWorkSize[1]; - program = - makePrintfProgram(&kernel, context, testId, testNum, formatNum, - isLongSupport, is64bAddressSpace(device)); + program = makePrintfProgram(&kernel, context, device, testId, + testNum, formatNum); if (!program || !kernel) { subtest_fail(nullptr); @@ -712,7 +869,12 @@ static int doTest(cl_command_queue queue, cl_context context, != verifyOutputBuffer(_analysisBuffer, allTestCase[testId], testNum, (cl_ulong)out32)) { - subtest_fail("verifyOutputBuffer failed\n"); + subtest_fail( + "verifyOutputBuffer failed with kernel: " + "\n%s\n expected: %s\n got: %s\n", + gLatestKernelSource.c_str(), + allTestCase[testId]->_correctBuffer[testNum].c_str(), + _analysisBuffer); continue; } } @@ -722,7 +884,12 @@ static int doTest(cl_command_queue queue, cl_context context, != verifyOutputBuffer(_analysisBuffer, allTestCase[testId], testNum, out64)) { - subtest_fail("verifyOutputBuffer failed\n"); + subtest_fail( + "verifyOutputBuffer failed with kernel: " + "\n%s\n expected: %s\n got: %s\n", + gLatestKernelSource.c_str(), + allTestCase[testId]->_correctBuffer[testNum].c_str(), + _analysisBuffer); continue; } } @@ -736,6 +903,8 @@ static int doTest(cl_command_queue queue, cl_context context, return s_test_fail - fail_count; } +} + int test_int(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) { @@ -814,6 +983,12 @@ int test_address_space(cl_device_id deviceID, cl_context context, return doTest(gQueue, gContext, TYPE_ADDRESS_SPACE, deviceID); } +int test_mixed_format_random(cl_device_id deviceID, cl_context context, + cl_command_queue queue, int num_elements) +{ + return doTest(gQueue, gContext, TYPE_MIXED_FORMAT_RANDOM, deviceID); +} + int test_buffer_size(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) { @@ -840,15 +1015,39 @@ int test_buffer_size(cl_device_id deviceID, cl_context context, } test_definition test_list[] = { - ADD_TEST(int), ADD_TEST(half), ADD_TEST(half_limits), - ADD_TEST(float), ADD_TEST(float_limits), ADD_TEST(octal), - ADD_TEST(unsigned), ADD_TEST(hexadecimal), ADD_TEST(char), - ADD_TEST(string), ADD_TEST(format_string), ADD_TEST(vector), - ADD_TEST(address_space), ADD_TEST(buffer_size), + ADD_TEST(int), + ADD_TEST(half), + ADD_TEST(half_limits), + ADD_TEST(float), + ADD_TEST(float_limits), + ADD_TEST(octal), + ADD_TEST(unsigned), + ADD_TEST(hexadecimal), + ADD_TEST(char), + ADD_TEST(string), + ADD_TEST(format_string), + ADD_TEST(vector), + ADD_TEST(address_space), + ADD_TEST(buffer_size), + ADD_TEST(mixed_format_random), }; const int test_num = ARRAY_SIZE( test_list ); +//----------------------------------------- +// printUsage +//----------------------------------------- +static void printUsage(void) +{ + log_info("test_printf: \n"); + log_info("\tdefault is to run the full test on the default device\n"); + log_info("\n"); + for (int i = 0; i < test_num; i++) + { + log_info("\t%s\n", test_list[i].name); + } +} + //----------------------------------------- // main //----------------------------------------- @@ -913,6 +1112,8 @@ int main(int argc, const char* argv[]) return -1; } + gMTdata = MTdataHolder(gRandomSeed); + int err = runTestHarnessWithCheck( argCount, argList, test_num, test_list, true, 0, InitCL ); if(gQueue) @@ -934,20 +1135,6 @@ int main(int argc, const char* argv[]) return err; } -//----------------------------------------- -// printUsage -//----------------------------------------- -static void printUsage( void ) -{ - log_info("test_printf: \n"); - log_info("\tdefault is to run the full test on the default device\n"); - log_info("\n"); - for( int i = 0; i < test_num; i++ ) - { - log_info( "\t%s\n", test_list[i].name ); - } -} - test_status InitCL( cl_device_id device ) { uint32_t device_frequency = 0; diff --git a/test_conformance/printf/test_printf.h b/test_conformance/printf/test_printf.h index 0a33d5f84a..a2cd9ed2be 100644 --- a/test_conformance/printf/test_printf.h +++ b/test_conformance/printf/test_printf.h @@ -58,6 +58,7 @@ enum PrintfTestType TYPE_FORMAT_STRING, TYPE_VECTOR, TYPE_ADDRESS_SPACE, + TYPE_MIXED_FORMAT_RANDOM, TYPE_COUNT }; diff --git a/test_conformance/printf/util_printf.cpp b/test_conformance/printf/util_printf.cpp index 29b7f8dc04..6e44b43fd8 100644 --- a/test_conformance/printf/util_printf.cpp +++ b/test_conformance/printf/util_printf.cpp @@ -1094,7 +1094,26 @@ testCase testCaseAddrSpace = { }; +//========================================================= +// mixed format +//========================================================= + +//---------------------------------------------------------- +// Container related to mixed format tests. +// Empty records for which the format string and reference string are generated +// at run time. The size of this vector specifies the number of random tests +// that will be run. +std::vector printMixedFormatGenParameters(64, + { { "" } }); +std::vector correctBufferMixedFormat; + +//---------------------------------------------------------- +// Test case for mixed-args +//---------------------------------------------------------- +testCase testCaseMixedFormat = { TYPE_MIXED_FORMAT_RANDOM, + correctBufferMixedFormat, + printMixedFormatGenParameters, NULL }; //------------------------------------------------------------------------------- @@ -1103,11 +1122,11 @@ testCase testCaseAddrSpace = { //------------------------------------------------------------------------------- std::vector allTestCase = { - &testCaseInt, &testCaseHalf, &testCaseHalfLimits, - &testCaseFloat, &testCaseFloatLimits, &testCaseOctal, - &testCaseUnsigned, &testCaseHexadecimal, &testCaseChar, - &testCaseString, &testCaseFormatString, &testCaseVector, - &testCaseAddrSpace + &testCaseInt, &testCaseHalf, &testCaseHalfLimits, + &testCaseFloat, &testCaseFloatLimits, &testCaseOctal, + &testCaseUnsigned, &testCaseHexadecimal, &testCaseChar, + &testCaseString, &testCaseFormatString, &testCaseVector, + &testCaseAddrSpace, &testCaseMixedFormat }; //----------------------------------------- @@ -1150,14 +1169,29 @@ size_t verifyOutputBuffer(char *analysisBuffer,testCase* pTestCase,size_t testId } - char* exp; - //Exponenent representation - if((exp = strstr(analysisBuffer,"E+")) != NULL || (exp = strstr(analysisBuffer,"e+")) != NULL || (exp = strstr(analysisBuffer,"E-")) != NULL || (exp = strstr(analysisBuffer,"e-")) != NULL) + char* exp = nullptr; + std::string copy_str; + std::vector staging(strlen(analysisBuffer) + 1); + std::vector staging_correct(pTestCase->_correctBuffer[testId].size() + + 1); + std::snprintf(staging.data(), staging.size(), "%s", analysisBuffer); + std::snprintf(staging_correct.data(), staging_correct.size(), "%s", + pTestCase->_correctBuffer[testId].c_str()); + // Exponenent representation + while ((exp = strstr(staging.data(), "E+")) != NULL + || (exp = strstr(staging.data(), "e+")) != NULL + || (exp = strstr(staging.data(), "E-")) != NULL + || (exp = strstr(staging.data(), "e-")) != NULL) { char correctExp[3]={0}; strncpy(correctExp,exp,2); - char* eCorrectBuffer = strstr((char*)pTestCase->_correctBuffer[testId].c_str(),correctExp); + // check if leading data is equal + int ret = strncmp(staging_correct.data(), staging.data(), + exp - staging.data()); + if (ret) return ret; + + char* eCorrectBuffer = strstr(staging_correct.data(), correctExp); if(eCorrectBuffer == NULL) return 1; @@ -1172,7 +1206,21 @@ size_t verifyOutputBuffer(char *analysisBuffer,testCase* pTestCase,size_t testId ++exp; while(*eCorrectBuffer == '0') ++eCorrectBuffer; - return strcmp(eCorrectBuffer,exp); + + copy_str = std::string(eCorrectBuffer); + std::snprintf(staging_correct.data(), staging_correct.size(), "%s", + copy_str.c_str()); + + copy_str = std::string(exp); + std::snprintf(staging.data(), staging.size(), "%s", copy_str.c_str()); + + if (strstr(staging.data(), "E+") != NULL + || strstr(staging.data(), "e+") != NULL + || strstr(staging.data(), "E-") != NULL + || strstr(staging.data(), "e-") != NULL) + continue; + + return strcmp(staging_correct.data(), copy_str.c_str()); } if (pTestCase->_correctBuffer[testId] == "inf") diff --git a/test_conformance/vulkan/CMakeLists.txt b/test_conformance/vulkan/CMakeLists.txt index 33eacffe2c..c970a77e38 100644 --- a/test_conformance/vulkan/CMakeLists.txt +++ b/test_conformance/vulkan/CMakeLists.txt @@ -1,10 +1,6 @@ set (MODULE_NAME VULKAN) -if(WIN32) - list(APPEND CLConform_LIBRARIES vulkan-1 vulkan_wrapper) -else(WIN32) - list(APPEND CLConform_LIBRARIES vulkan dl vulkan_wrapper) -endif(WIN32) +list(APPEND CLConform_LIBRARIES vulkan_wrapper) set(CMAKE_CXX_FLAGS "-fpermissive") if(WIN32) set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -DVK_USE_PLATFORM_WIN32_KHR")