From 3349709dd0647b42c69851805d32ad6caa743968 Mon Sep 17 00:00:00 2001 From: Marcin Hajder Date: Thu, 18 Apr 2024 16:43:32 +0200 Subject: [PATCH 1/7] Extended printf test with literal strings to verify special symbols and format string -correction to remove duplication of separate tests for each element of vector -added new sort of test to verify format string -extended previous string tests with special and nested symbols --- test_conformance/printf/test_printf.cpp | 889 ++++++++++-------------- test_conformance/printf/test_printf.h | 3 +- test_conformance/printf/util_printf.cpp | 459 ++++++++---- 3 files changed, 689 insertions(+), 662 deletions(-) diff --git a/test_conformance/printf/test_printf.cpp b/test_conformance/printf/test_printf.cpp index e43e302f1..0c5eb169a 100644 --- a/test_conformance/printf/test_printf.cpp +++ b/test_conformance/printf/test_printf.cpp @@ -14,6 +14,7 @@ // limitations under the License. // #include "harness/os_helpers.h" +#include "harness/typeWrappers.h" #include #include @@ -78,10 +79,15 @@ static 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,bool isLongSupport = true,bool is64bAddrSpace = false); +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); // 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, const unsigned int testNum, cl_device_id device); +static 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); @@ -102,6 +108,7 @@ int waitForEvent(cl_event* event); // tracks the subtests int s_test_cnt = 0; int s_test_fail = 0; +int s_test_skip = 0; static cl_context gContext; @@ -159,7 +166,9 @@ static void getAnalysisBuffer(char* analysisBuffer) if(NULL == fp) log_error("Failed to open analysis buffer ('%s')\n", strerror(errno)); else - while(fgets(analysisBuffer, ANALYSIS_BUFFER_SIZE, fp) != NULL ); + std::fread(analysisBuffer, sizeof(analysisBuffer[0]), + ANALYSIS_BUFFER_SIZE, fp); + fclose(fp); } @@ -206,7 +215,12 @@ int waitForEvent(cl_event* event) //----------------------------------------- // makePrintfProgram //----------------------------------------- -static cl_program makePrintfProgram(cl_kernel *kernel_ptr, const cl_context context,const unsigned int testId,const unsigned int testNum,bool isLongSupport,bool is64bAddrSpace) +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) { int err; cl_program program; @@ -215,20 +229,6 @@ static cl_program makePrintfProgram(cl_kernel *kernel_ptr, const cl_context cont char addrSpacePAddArgument[256] = {0}; char extension[128] = { 0 }; - //Program Source code for int,float,octal,hexadecimal,char,string - const char* sourceGen[] = { - extension, - "__kernel void ", - testname, - "(void)\n", - "{\n" - " printf(\"", - allTestCase[testId]->_genParameters[testNum].genericFormat, - "\\n\",", - allTestCase[testId]->_genParameters[testNum].dataRepresentation, - ");", - "}\n" - }; //Program Source code for vector const char* sourceVec[] = { extension, @@ -254,12 +254,20 @@ static cl_program makePrintfProgram(cl_kernel *kernel_ptr, const cl_context cont "}\n" }; //Program Source code for address space - const char *sourceAddrSpace[] = { - "__kernel void ", testname,"(",addrSpaceArgument, + const char* sourceAddrSpace[] = { + "__kernel void ", + testname, + "(", + addrSpaceArgument, ")\n{\n", - allTestCase[testId]->_genParameters[testNum].addrSpaceVariableTypeQualifier, + allTestCase[testId] + ->_genParameters[testNum] + .addrSpaceVariableTypeQualifier, "printf(", - allTestCase[testId]->_genParameters[testNum].genericFormat, + allTestCase[testId] + ->_genParameters[testNum] + .genericFormats[formatNum] + .c_str(), ",", allTestCase[testId]->_genParameters[testNum].addrSpaceParameter, "); ", @@ -314,9 +322,33 @@ static cl_program makePrintfProgram(cl_kernel *kernel_ptr, const cl_context cont } else { - err = create_single_kernel_helper( - context, &program, kernel_ptr, - sizeof(sourceGen) / sizeof(sourceGen[0]), sourceGen, testname); + // Program Source code for int,float,octal,hexadecimal,char,string + std::ostringstream sourceGen; + sourceGen << extension << "__kernel void " << testname + << "(void)\n" + "{\n" + " printf(\"" + << allTestCase[testId] + ->_genParameters[testNum] + .genericFormats[formatNum] + .c_str() + << "\\n\""; + + if (allTestCase[testId]->_genParameters[testNum].dataRepresentation) + { + sourceGen << "," + << allTestCase[testId] + ->_genParameters[testNum] + .dataRepresentation; + } + + sourceGen << ");}\n"; + + std::string kernel_source = sourceGen.str(); + const char* ptr = kernel_source.c_str(); + + err = create_single_kernel_helper(context, &program, kernel_ptr, 1, + &ptr, testname); } if (!program || err) { @@ -396,521 +428,377 @@ static bool is64bAddressSpace(cl_device_id device_id) else return false; } + +//----------------------------------------- +// subtest_fail +//----------------------------------------- +template void subtest_fail(const char* msg, Args... args) +{ + if (msg) log_error(msg, args...); + ++s_test_fail; + ++s_test_cnt; +}; + //----------------------------------------- // doTest //----------------------------------------- -static int doTest(cl_command_queue queue, cl_context context, const unsigned int testId, const unsigned int testNum, cl_device_id device) +static int doTest(cl_command_queue queue, cl_context context, + const unsigned int testId, cl_device_id device) { + int err = CL_SUCCESS; + if ((allTestCase[testId]->_type == TYPE_HALF || allTestCase[testId]->_type == TYPE_HALF_LIMITS) && !is_extension_available(device, "cl_khr_fp16")) { - log_info( - "Skipping half because cl_khr_fp16 extension is not supported.\n"); + log_info("Skipping half because cl_khr_fp16 extension is not " + "supported.\n"); return TEST_SKIPPED_ITSELF; } - if(allTestCase[testId]->_type == TYPE_VECTOR) - { - if ((strcmp(allTestCase[testId]->_genParameters[testNum].dataType, - "half") - == 0) - && !is_extension_available(device, "cl_khr_fp16")) - { - log_info("Skipping half because cl_khr_fp16 extension is not " - "supported.\n"); - return TEST_SKIPPED_ITSELF; - } + auto& genParams = allTestCase[testId]->_genParameters; - 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) + auto fail_count = s_test_fail; + auto pass_count = s_test_cnt; + auto skip_count = s_test_skip; + + for (unsigned testNum = 0; testNum < genParams.size(); testNum++) { - if(isKernelArgument(allTestCase[testId], testNum)) + for (unsigned formatNum = 0; formatNum < allTestCase[testId] + ->_genParameters[testNum] + .genericFormats.size(); + formatNum++) { - log_info("%d)testing kernel //argument %s \n printf(%s,%s)\n",testNum,allTestCase[testId]->_genParameters[testNum].addrSpaceArgumentTypeQualifier, - allTestCase[testId]->_genParameters[testNum].genericFormat,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].genericFormat,allTestCase[testId]->_genParameters[testNum].addrSpaceParameter); - } - } - else - { - log_info("%d)testing printf(\"%s\",%s)\n",testNum,allTestCase[testId]->_genParameters[testNum].genericFormat,allTestCase[testId]->_genParameters[testNum].dataRepresentation); - } + if (allTestCase[testId]->_type == TYPE_VECTOR) + { + if ((strcmp( + allTestCase[testId]->_genParameters[testNum].dataType, + "half") + == 0) + && !is_extension_available(device, "cl_khr_fp16")) + { + log_info( + "Skipping half because cl_khr_fp16 extension is not " + "supported.\n"); + + s_test_skip++; + s_test_cnt++; + continue; + } - // Long support for varible type - if(allTestCase[testId]->_type == TYPE_VECTOR && !strcmp(allTestCase[testId]->_genParameters[testNum].dataType,"long") && !isLongSupported(device)) - { - log_info( "Long is not supported, test not run.\n" ); - return 0; - } + 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\",%s)\n", testNum, + allTestCase[testId] + ->_genParameters[testNum] + .genericFormats[formatNum] + .c_str(), + allTestCase[testId] + ->_genParameters[testNum] + .dataRepresentation); + } - // 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; - } + // Long support for varible type + if (allTestCase[testId]->_type == TYPE_VECTOR + && !strcmp( + allTestCase[testId]->_genParameters[testNum].dataType, + "long") + && !isLongSupported(device)) + { + log_info("Long is not supported, test not run.\n"); + s_test_skip++; + s_test_cnt++; + continue; + } - int err; - cl_program program; - cl_kernel kernel; - cl_mem d_out = NULL; - cl_mem d_a = NULL; - char _analysisBuffer[ANALYSIS_BUFFER_SIZE]; - cl_uint out32 = 0; - cl_ulong out64 = 0; - int fd = -1; - - // Define an index space (global work size) of threads for execution. - size_t globalWorkSize[1]; - - program = makePrintfProgram(&kernel, context,testId,testNum,isLongSupport,is64bAddressSpace(device)); - if (!program || !kernel) { - ++s_test_fail; - ++s_test_cnt; - return -1; - } + fflush(stdout); - //For address space test if there is kernel argument - set it - if(allTestCase[testId]->_type == TYPE_ADDRESS_SPACE ) - { - if(isKernelArgument(allTestCase[testId],testNum)) - { - int a = 2; - d_a = clCreateBuffer(context, CL_MEM_READ_ONLY|CL_MEM_COPY_HOST_PTR, - sizeof(int), &a, &err); - if(err!= CL_SUCCESS || d_a == NULL) { - log_error("clCreateBuffer failed\n"); - goto exit; + // 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; } - err = clSetKernelArg(kernel, 0, sizeof(cl_mem), &d_a); - if(err!= CL_SUCCESS) { - log_error("clSetKernelArg failed\n"); - goto exit; + + clProgramWrapper program; + clKernelWrapper kernel; + clMemWrapper d_out; + clMemWrapper d_a; + char _analysisBuffer[ANALYSIS_BUFFER_SIZE]; + cl_uint out32 = 0; + cl_ulong out64 = 0; + int fd = -1; + + // Define an index space (global work size) of threads for + // execution. + size_t globalWorkSize[1]; + + program = + makePrintfProgram(&kernel, context, testId, testNum, formatNum, + isLongSupport, is64bAddressSpace(device)); + if (!program || !kernel) + { + subtest_fail(nullptr); + continue; } - } - //For address space test if %p is tested - if(isKernelPFormat(allTestCase[testId],testNum)) - { - d_out = clCreateBuffer(context, CL_MEM_READ_WRITE, - sizeof(cl_ulong), NULL, &err); - if(err!= CL_SUCCESS || d_out == NULL) { - log_error("clCreateBuffer failed\n"); - goto exit; + + // For address space test if there is kernel argument - set it + if (allTestCase[testId]->_type == TYPE_ADDRESS_SPACE) + { + if (isKernelArgument(allTestCase[testId], testNum)) + { + int a = 2; + d_a = clCreateBuffer( + context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, + sizeof(int), &a, &err); + if (err != CL_SUCCESS || d_a == NULL) + { + subtest_fail("clCreateBuffer failed\n"); + continue; + } + err = clSetKernelArg(kernel, 0, sizeof(cl_mem), &d_a); + if (err != CL_SUCCESS) + { + subtest_fail("clSetKernelArg failed\n"); + continue; + } + } + // For address space test if %p is tested + if (isKernelPFormat(allTestCase[testId], testNum)) + { + d_out = clCreateBuffer(context, CL_MEM_READ_WRITE, + sizeof(cl_ulong), NULL, &err); + if (err != CL_SUCCESS || d_out == NULL) + { + subtest_fail("clCreateBuffer failed\n"); + continue; + } + err = clSetKernelArg(kernel, 1, sizeof(cl_mem), &d_out); + if (err != CL_SUCCESS) + { + subtest_fail("clSetKernelArg failed\n"); + continue; + } + } } - err = clSetKernelArg(kernel, 1, sizeof(cl_mem), &d_out); - if(err!= CL_SUCCESS) { - log_error("clSetKernelArg failed\n"); - goto exit; + + fd = acquireOutputStream(&err); + if (err != 0) + { + subtest_fail("Error while redirection stdout to file"); + continue; + } + globalWorkSize[0] = 1; + cl_event ndrEvt; + err = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, globalWorkSize, + NULL, 0, NULL, &ndrEvt); + if (err != CL_SUCCESS) + { + releaseOutputStream(fd); + subtest_fail("\n clEnqueueNDRangeKernel failed errcode:%d\n", + err); + continue; } - } - } - fd = acquireOutputStream(&err); - if (err != 0) - { - log_error("Error while redirection stdout to file"); - goto exit; - } - globalWorkSize[0] = 1; - cl_event ndrEvt; - err = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, globalWorkSize, NULL, 0, NULL,&ndrEvt); - if (err != CL_SUCCESS) { - releaseOutputStream(fd); - log_error("\n clEnqueueNDRangeKernel failed errcode:%d\n", err); - ++s_test_fail; - goto exit; - } + fflush(stdout); + err = clFlush(queue); + if (err != CL_SUCCESS) + { + releaseOutputStream(fd); + subtest_fail("clFlush failed : %d\n", err); + continue; + } + // Wait until kernel finishes its execution and (thus) the output + // printed from the kernel is immediately printed + err = waitForEvent(&ndrEvt); - fflush(stdout); - err = clFlush(queue); - if(err != CL_SUCCESS) - { - releaseOutputStream(fd); - log_error("clFlush failed\n"); - goto exit; - } - //Wait until kernel finishes its execution and (thus) the output printed from the kernel - //is immediately printed - err = waitForEvent(&ndrEvt); + releaseOutputStream(fd); - releaseOutputStream(fd); + if (err != CL_SUCCESS) + { + subtest_fail("waitforEvent failed : %d\n", err); + continue; + } + fflush(stdout); - if(err != CL_SUCCESS) - { - log_error("waitforEvent failed\n"); - goto exit; - } - fflush(stdout); + if (allTestCase[testId]->_type == TYPE_ADDRESS_SPACE + && isKernelPFormat(allTestCase[testId], testNum)) + { + // Read the OpenCL output buffer (d_out) to the host output + // array (out) + if (!is64bAddressSpace(device)) // 32-bit address space + { + clEnqueueReadBuffer(queue, d_out, CL_TRUE, 0, + sizeof(cl_int), &out32, 0, NULL, NULL); + } + else // 64-bit address space + { + clEnqueueReadBuffer(queue, d_out, CL_TRUE, 0, + sizeof(cl_ulong), &out64, 0, NULL, + NULL); + } + } - if(allTestCase[testId]->_type == TYPE_ADDRESS_SPACE && isKernelPFormat(allTestCase[testId],testNum)) - { - // Read the OpenCL output buffer (d_out) to the host output array (out) - if(!is64bAddressSpace(device))//32-bit address space - { - clEnqueueReadBuffer(queue, d_out, CL_TRUE, 0, sizeof(cl_int),&out32, - 0, NULL, NULL); - } - else //64-bit address space - { - clEnqueueReadBuffer(queue, d_out, CL_TRUE, 0, sizeof(cl_ulong),&out64, - 0, NULL, NULL); + // + // Get the output printed from the kernel to _analysisBuffer + // and verify its correctness + getAnalysisBuffer(_analysisBuffer); + if (!is64bAddressSpace(device)) // 32-bit address space + { + if (0 + != verifyOutputBuffer(_analysisBuffer, allTestCase[testId], + testNum, (cl_ulong)out32)) + { + subtest_fail("verifyOutputBuffer failed\n"); + continue; + } + } + else // 64-bit address space + { + if (0 + != verifyOutputBuffer(_analysisBuffer, allTestCase[testId], + testNum, out64)) + { + subtest_fail("verifyOutputBuffer failed\n"); + continue; + } + } + ++s_test_cnt; } } - // - //Get the output printed from the kernel to _analysisBuffer - //and verify its correctness - getAnalysisBuffer(_analysisBuffer); - if(!is64bAddressSpace(device)) //32-bit address space - { - if(0 != verifyOutputBuffer(_analysisBuffer,allTestCase[testId],testNum,(cl_ulong) out32)) - err = ++s_test_fail; - } - else //64-bit address space - { - if(0 != verifyOutputBuffer(_analysisBuffer,allTestCase[testId],testNum,out64)) - err = ++s_test_fail; - } -exit: - if(clReleaseKernel(kernel) != CL_SUCCESS) - log_error("clReleaseKernel failed\n"); - if(clReleaseProgram(program) != CL_SUCCESS) - log_error("clReleaseProgram failed\n"); - if(d_out) - clReleaseMemObject(d_out); - if(d_a) - clReleaseMemObject(d_a); - ++s_test_cnt; - return err; + // all subtests skipped ? + if (s_test_skip - skip_count == s_test_cnt - pass_count) + return TEST_SKIPPED_ITSELF; + return s_test_fail - fail_count; } - -int test_int_0(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) -{ - return doTest(gQueue, gContext, TYPE_INT, 0, deviceID); -} -int test_int_1(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) -{ - return doTest(gQueue, gContext, TYPE_INT, 1, deviceID); -} -int test_int_2(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) -{ - return doTest(gQueue, gContext, TYPE_INT, 2, deviceID); -} -int test_int_3(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) -{ - return doTest(gQueue, gContext, TYPE_INT, 3, deviceID); -} -int test_int_4(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) -{ - return doTest(gQueue, gContext, TYPE_INT, 4, deviceID); -} -int test_int_5(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) -{ - return doTest(gQueue, gContext, TYPE_INT, 5, deviceID); -} -int test_int_6(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) -{ - return doTest(gQueue, gContext, TYPE_INT, 6, deviceID); -} -int test_int_7(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) -{ - return doTest(gQueue, gContext, TYPE_INT, 7, deviceID); -} -int test_int_8(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) +int test_int(cl_device_id deviceID, cl_context context, cl_command_queue queue, + int num_elements) { - return doTest(gQueue, gContext, TYPE_INT, 8, deviceID); + return doTest(gQueue, gContext, TYPE_INT, deviceID); } - -int test_half_0(cl_device_id deviceID, cl_context context, - cl_command_queue queue, int num_elements) -{ - return doTest(gQueue, gContext, TYPE_HALF, 0, deviceID); -} -int test_half_1(cl_device_id deviceID, cl_context context, - cl_command_queue queue, int num_elements) -{ - return doTest(gQueue, gContext, TYPE_HALF, 1, deviceID); -} -int test_half_2(cl_device_id deviceID, cl_context context, - cl_command_queue queue, int num_elements) -{ - return doTest(gQueue, gContext, TYPE_HALF, 2, deviceID); -} -int test_half_3(cl_device_id deviceID, cl_context context, - cl_command_queue queue, int num_elements) -{ - return doTest(gQueue, gContext, TYPE_HALF, 3, deviceID); -} -int test_half_4(cl_device_id deviceID, cl_context context, - cl_command_queue queue, int num_elements) -{ - return doTest(gQueue, gContext, TYPE_HALF, 4, deviceID); -} -int test_half_5(cl_device_id deviceID, cl_context context, - cl_command_queue queue, int num_elements) +int test_half(cl_device_id deviceID, cl_context context, cl_command_queue queue, + int num_elements) { - return doTest(gQueue, gContext, TYPE_HALF, 5, deviceID); + return doTest(gQueue, gContext, TYPE_HALF, deviceID); } -int test_half_6(cl_device_id deviceID, cl_context context, - cl_command_queue queue, int num_elements) -{ - return doTest(gQueue, gContext, TYPE_HALF, 6, deviceID); -} -int test_half_7(cl_device_id deviceID, cl_context context, - cl_command_queue queue, int num_elements) -{ - return doTest(gQueue, gContext, TYPE_HALF, 7, deviceID); -} -int test_half_8(cl_device_id deviceID, cl_context context, - cl_command_queue queue, int num_elements) -{ - return doTest(gQueue, gContext, TYPE_HALF, 8, deviceID); -} -int test_half_9(cl_device_id deviceID, cl_context context, - cl_command_queue queue, int num_elements) -{ - return doTest(gQueue, gContext, TYPE_HALF, 9, deviceID); -} - -int test_half_limits_0(cl_device_id deviceID, cl_context context, - cl_command_queue queue, int num_elements) -{ - return doTest(gQueue, gContext, TYPE_HALF_LIMITS, 0, deviceID); -} -int test_half_limits_1(cl_device_id deviceID, cl_context context, - cl_command_queue queue, int num_elements) -{ - return doTest(gQueue, gContext, TYPE_HALF_LIMITS, 1, deviceID); -} -int test_half_limits_2(cl_device_id deviceID, cl_context context, - cl_command_queue queue, int num_elements) +int test_half_limits(cl_device_id deviceID, cl_context context, + cl_command_queue queue, int num_elements) { - return doTest(gQueue, gContext, TYPE_HALF_LIMITS, 2, deviceID); + return doTest(gQueue, gContext, TYPE_HALF_LIMITS, deviceID); } - -int test_float_0(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) -{ - return doTest(gQueue, gContext, TYPE_FLOAT, 0, deviceID); -} -int test_float_1(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) -{ - return doTest(gQueue, gContext, TYPE_FLOAT, 1, deviceID); -} -int test_float_2(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) -{ - return doTest(gQueue, gContext, TYPE_FLOAT, 2, deviceID); -} -int test_float_3(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) -{ - return doTest(gQueue, gContext, TYPE_FLOAT, 3, deviceID); -} -int test_float_4(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) -{ - return doTest(gQueue, gContext, TYPE_FLOAT, 4, deviceID); -} -int test_float_5(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) -{ - return doTest(gQueue, gContext, TYPE_FLOAT, 5, deviceID); -} -int test_float_6(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) -{ - return doTest(gQueue, gContext, TYPE_FLOAT, 6, deviceID); -} -int test_float_7(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) -{ - return doTest(gQueue, gContext, TYPE_FLOAT, 7, deviceID); -} -int test_float_8(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) -{ - return doTest(gQueue, gContext, TYPE_FLOAT, 8, deviceID); -} -int test_float_9(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) -{ - return doTest(gQueue, gContext, TYPE_FLOAT, 9, deviceID); -} -int test_float_10(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) -{ - return doTest(gQueue, gContext, TYPE_FLOAT, 10, deviceID); -} -int test_float_11(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) -{ - return doTest(gQueue, gContext, TYPE_FLOAT, 11, deviceID); -} -int test_float_12(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) -{ - return doTest(gQueue, gContext, TYPE_FLOAT, 12, deviceID); -} -int test_float_13(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) -{ - return doTest(gQueue, gContext, TYPE_FLOAT, 13, deviceID); -} -int test_float_14(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) +int test_float(cl_device_id deviceID, cl_context context, + cl_command_queue queue, int num_elements) { - return doTest(gQueue, gContext, TYPE_FLOAT, 14, deviceID); + return doTest(gQueue, gContext, TYPE_FLOAT, deviceID); } -int test_float_15(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) -{ - return doTest(gQueue, gContext, TYPE_FLOAT, 15, deviceID); -} -int test_float_16(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) -{ - return doTest(gQueue, gContext, TYPE_FLOAT, 16, deviceID); -} -int test_float_17(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) -{ - return doTest(gQueue, gContext, TYPE_FLOAT, 17, deviceID); -} - -int test_float_limits_0(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) -{ - return doTest(gQueue, gContext, TYPE_FLOAT_LIMITS, 0, deviceID); -} -int test_float_limits_1(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) +int test_float_limits(cl_device_id deviceID, cl_context context, + cl_command_queue queue, int num_elements) { - return doTest(gQueue, gContext, TYPE_FLOAT_LIMITS, 1, deviceID); -} -int test_float_limits_2(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) -{ - return doTest(gQueue, gContext, TYPE_FLOAT_LIMITS, 2, deviceID); + return doTest(gQueue, gContext, TYPE_FLOAT_LIMITS, deviceID); } - -int test_octal_0(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) -{ - return doTest(gQueue, gContext, TYPE_OCTAL, 0, deviceID); -} -int test_octal_1(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) -{ - return doTest(gQueue, gContext, TYPE_OCTAL, 1, deviceID); -} -int test_octal_2(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) +int test_octal(cl_device_id deviceID, cl_context context, + cl_command_queue queue, int num_elements) { - return doTest(gQueue, gContext, TYPE_OCTAL, 2, deviceID); + return doTest(gQueue, gContext, TYPE_OCTAL, deviceID); } -int test_octal_3(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) -{ - return doTest(gQueue, gContext, TYPE_OCTAL, 3, deviceID); -} - -int test_unsigned_0(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) -{ - return doTest(gQueue, gContext, TYPE_UNSIGNED, 0, deviceID); -} -int test_unsigned_1(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) +int test_unsigned(cl_device_id deviceID, cl_context context, + cl_command_queue queue, int num_elements) { - return doTest(gQueue, gContext, TYPE_UNSIGNED, 1, deviceID); + return doTest(gQueue, gContext, TYPE_UNSIGNED, deviceID); } - -int test_hexadecimal_0(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) -{ - return doTest(gQueue, gContext, TYPE_HEXADEC, 0, deviceID); -} -int test_hexadecimal_1(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) -{ - return doTest(gQueue, gContext, TYPE_HEXADEC, 1, deviceID); -} -int test_hexadecimal_2(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) -{ - return doTest(gQueue, gContext, TYPE_HEXADEC, 2, deviceID); -} -int test_hexadecimal_3(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) -{ - return doTest(gQueue, gContext, TYPE_HEXADEC, 3, deviceID); -} -int test_hexadecimal_4(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) +int test_hexadecimal(cl_device_id deviceID, cl_context context, + cl_command_queue queue, int num_elements) { - return doTest(gQueue, gContext, TYPE_HEXADEC, 4, deviceID); + return doTest(gQueue, gContext, TYPE_HEXADEC, deviceID); } - -int test_char_0(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) +int test_char(cl_device_id deviceID, cl_context context, cl_command_queue queue, + int num_elements) { - return doTest(gQueue, gContext, TYPE_CHAR, 0, deviceID); + return doTest(gQueue, gContext, TYPE_CHAR, deviceID); } -int test_char_1(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) -{ - return doTest(gQueue, gContext, TYPE_CHAR, 1, deviceID); -} -int test_char_2(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) -{ - return doTest(gQueue, gContext, TYPE_CHAR, 2, deviceID); -} - -int test_string_0(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) -{ - return doTest(gQueue, gContext, TYPE_STRING, 0, deviceID); -} -int test_string_1(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) -{ - return doTest(gQueue, gContext, TYPE_STRING, 1, deviceID); -} -int test_string_2(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) +int test_string(cl_device_id deviceID, cl_context context, + cl_command_queue queue, int num_elements) { - return doTest(gQueue, gContext, TYPE_STRING, 2, deviceID); + return doTest(gQueue, gContext, TYPE_STRING, deviceID); } - -int test_vector_0(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) -{ - return doTest(gQueue, gContext, TYPE_VECTOR, 0, deviceID); -} -int test_vector_1(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) -{ - return doTest(gQueue, gContext, TYPE_VECTOR, 1, deviceID); -} -int test_vector_2(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) -{ - return doTest(gQueue, gContext, TYPE_VECTOR, 2, deviceID); -} -int test_vector_3(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) -{ - return doTest(gQueue, gContext, TYPE_VECTOR, 3, deviceID); -} -int test_vector_4(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) -{ - return doTest(gQueue, gContext, TYPE_VECTOR, 4, deviceID); -} -int test_vector_5(cl_device_id deviceID, cl_context context, - cl_command_queue queue, int num_elements) +int test_format_string(cl_device_id deviceID, cl_context context, + cl_command_queue queue, int num_elements) { - return doTest(gQueue, gContext, TYPE_VECTOR, 5, deviceID); + return doTest(gQueue, gContext, TYPE_FORMAT_STRING, deviceID); } - -int test_address_space_0(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) -{ - return doTest(gQueue, gContext, TYPE_ADDRESS_SPACE, 0, deviceID); -} -int test_address_space_1(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) -{ - return doTest(gQueue, gContext, TYPE_ADDRESS_SPACE, 1, deviceID); -} -int test_address_space_2(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) -{ - return doTest(gQueue, gContext, TYPE_ADDRESS_SPACE, 2, deviceID); -} -int test_address_space_3(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) +int test_vector(cl_device_id deviceID, cl_context context, + cl_command_queue queue, int num_elements) { - return doTest(gQueue, gContext, TYPE_ADDRESS_SPACE, 3, deviceID); + return doTest(gQueue, gContext, TYPE_VECTOR, deviceID); } -int test_address_space_4(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) + +int test_address_space(cl_device_id deviceID, cl_context context, + cl_command_queue queue, int num_elements) { - return doTest(gQueue, gContext, TYPE_ADDRESS_SPACE, 4, deviceID); + return doTest(gQueue, gContext, TYPE_ADDRESS_SPACE, deviceID); } int test_buffer_size(cl_device_id deviceID, cl_context context, @@ -939,58 +827,11 @@ int test_buffer_size(cl_device_id deviceID, cl_context context, } test_definition test_list[] = { - ADD_TEST(int_0), ADD_TEST(int_1), - ADD_TEST(int_2), ADD_TEST(int_3), - ADD_TEST(int_4), ADD_TEST(int_5), - ADD_TEST(int_6), ADD_TEST(int_7), - ADD_TEST(int_8), - - ADD_TEST(half_0), ADD_TEST(half_1), - ADD_TEST(half_2), ADD_TEST(half_3), - ADD_TEST(half_4), ADD_TEST(half_5), - ADD_TEST(half_6), ADD_TEST(half_7), - ADD_TEST(half_8), ADD_TEST(half_9), - - ADD_TEST(half_limits_0), ADD_TEST(half_limits_1), - ADD_TEST(half_limits_2), - - ADD_TEST(float_0), ADD_TEST(float_1), - ADD_TEST(float_2), ADD_TEST(float_3), - ADD_TEST(float_4), ADD_TEST(float_5), - ADD_TEST(float_6), ADD_TEST(float_7), - ADD_TEST(float_8), ADD_TEST(float_9), - ADD_TEST(float_10), ADD_TEST(float_11), - ADD_TEST(float_12), ADD_TEST(float_13), - ADD_TEST(float_14), ADD_TEST(float_15), - ADD_TEST(float_16), ADD_TEST(float_17), - - ADD_TEST(float_limits_0), ADD_TEST(float_limits_1), - ADD_TEST(float_limits_2), - - ADD_TEST(octal_0), ADD_TEST(octal_1), - ADD_TEST(octal_2), ADD_TEST(octal_3), - - ADD_TEST(unsigned_0), ADD_TEST(unsigned_1), - - ADD_TEST(hexadecimal_0), ADD_TEST(hexadecimal_1), - ADD_TEST(hexadecimal_2), ADD_TEST(hexadecimal_3), - ADD_TEST(hexadecimal_4), - - ADD_TEST(char_0), ADD_TEST(char_1), - ADD_TEST(char_2), - - ADD_TEST(string_0), ADD_TEST(string_1), - ADD_TEST(string_2), - - ADD_TEST(vector_0), ADD_TEST(vector_1), - ADD_TEST(vector_2), ADD_TEST(vector_3), - ADD_TEST(vector_4), ADD_TEST(vector_5), - - ADD_TEST(address_space_0), ADD_TEST(address_space_1), - ADD_TEST(address_space_2), ADD_TEST(address_space_3), - ADD_TEST(address_space_4), - - 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), }; const int test_num = ARRAY_SIZE( test_list ); diff --git a/test_conformance/printf/test_printf.h b/test_conformance/printf/test_printf.h index 8eb2a0324..0a33d5f84 100644 --- a/test_conformance/printf/test_printf.h +++ b/test_conformance/printf/test_printf.h @@ -55,6 +55,7 @@ enum PrintfTestType TYPE_HEXADEC, TYPE_CHAR, TYPE_STRING, + TYPE_FORMAT_STRING, TYPE_VECTOR, TYPE_ADDRESS_SPACE, TYPE_COUNT @@ -62,7 +63,7 @@ enum PrintfTestType struct printDataGenParameters { - const char* genericFormat; + std::vector genericFormats; const char* dataRepresentation; const char* vectorFormatFlag; const char* vectorFormatSpecifier; diff --git a/test_conformance/printf/util_printf.cpp b/test_conformance/printf/util_printf.cpp index 6b310a994..deb627958 100644 --- a/test_conformance/printf/util_printf.cpp +++ b/test_conformance/printf/util_printf.cpp @@ -46,39 +46,41 @@ std::vector printIntGenParameters = { //(Minimum)Five-wide,default(right)-justified - {"%5d","10"}, + { { "%5d" }, "10" }, - //(Minimum)Five-wide,left-justified + //(Minimum)Five-wide,left-justified - {"%-5d","10"}, + { { "%-5d" }, "10" }, - //(Minimum)Five-wide,default(right)-justified,zero-filled + //(Minimum)Five-wide,default(right)-justified,zero-filled - {"%05d","10"}, + { { "%05d" }, "10" }, - //(Minimum)Five-wide,default(right)-justified,with sign + //(Minimum)Five-wide,default(right)-justified,with sign - {"%+5d","10"}, + { { "%+5d" }, "10" }, - //(Minimum)Five-wide ,left-justified,with sign + //(Minimum)Five-wide ,left-justified,with sign - {"%-+5d","10"}, + { { "%-+5d" }, "10" }, - //(Minimum)Five-digit(zero-filled in absent digits),default(right)-justified + //(Minimum)Five-digit(zero-filled in absent digits),default(right)-justified - {"%.5i","100"}, + { { "%.5i" }, "100" }, - //(Minimum)Six-wide,Five-digit(zero-filled in absent digits),default(right)-justified + //(Minimum)Six-wide,Five-digit(zero-filled in absent + // digits),default(right)-justified - {"%6.5i","100"}, + { { "%6.5i" }, "100" }, - //0 and - flag both apper ==>0 is ignored,left-justified,capital I + // 0 and - flag both apper ==>0 is ignored,left-justified,capital I - {"%-06i","100"}, + { { "%-06i" }, "100" }, - //(Minimum)Six-wide,Five-digit(zero-filled in absent digits),default(right)-justified + //(Minimum)Six-wide,Five-digit(zero-filled in absent + // digits),default(right)-justified - {"%06.5i","100"} + { { "%06.5i" }, "100" } }; @@ -119,50 +121,50 @@ std::vector printHalfGenParameters = { // Default(right)-justified - { "%f", "1.234h" }, + { { "%f" }, "1.234h" }, // One position after the decimal,default(right)-justified - { "%4.2f", "1.2345h" }, + { { "%4.2f" }, "1.2345h" }, // Zero positions after the // decimal([floor]rounding),default(right)-justified - { "%.0f", "0.1h" }, + { { "%.0f" }, "0.1h" }, // Zero positions after the decimal([ceil]rounding),default(right)-justified - { "%.0f", "0.6h" }, + { { "%.0f" }, "0.6h" }, // Zero-filled,default positions number after the // decimal,default(right)-justified - { "%0f", "0.6h" }, + { { "%0f" }, "0.6h" }, // Double argument representing floating-point,used by f // style,default(right)-justified - { "%4g", "5.678h" }, + { { "%4g" }, "5.678h" }, // Double argument representing floating-point,used by e // style,default(right)-justified - { "%4.2g", "5.678h" }, + { { "%4.2g" }, "5.678h" }, // Double argument representing floating-point,used by e // style,default(right)-justified - { "%4G", "0.000062h" }, + { { "%4G" }, "0.000062h" }, // Double argument representing floating-point,with // exponent,left-justified,default(right)-justified - { "%-#20.15e", "65504.0h" }, + { { "%-#20.15e" }, "65504.0h" }, // Double argument representing floating-point,with // exponent,left-justified,with sign,capital E,default(right)-justified - { "%+#21.15E", "-65504.0h" }, + { { "%+#21.15E" }, "-65504.0h" }, }; //--------------------------------------------------------- @@ -204,14 +206,14 @@ std::vector printHalfLimitsGenParameters = { // Infinity (1.0/0.0) - { "%f", "1.0h/0.0h" }, + { { "%f" }, "1.0h/0.0h" }, // NaN - { "%f", "sqrt(-1.0h)" }, + { { "%f" }, "sqrt(-1.0h)" }, // NaN - { "%f", "acospi(2.0h)" } + { { "%f" }, "acospi(2.0h)" } }; //-------------------------------------------------------- @@ -265,77 +267,89 @@ testCase testCaseHalfLimits = { std::vector printFloatGenParameters = { - //Default(right)-justified + // Default(right)-justified - {"%f","10.3456"}, + { { "%f" }, "10.3456" }, - //One position after the decimal,default(right)-justified + // One position after the decimal,default(right)-justified - {"%.1f","10.3456"}, + { { "%.1f" }, "10.3456" }, - //Two positions after the decimal,default(right)-justified + // Two positions after the decimal,default(right)-justified - {"%.2f","10.3456"}, + { { "%.2f" }, "10.3456" }, - //(Minimum)Eight-wide,three positions after the decimal,default(right)-justified + //(Minimum)Eight-wide,three positions after the + // decimal,default(right)-justified - {"%8.3f","10.3456"}, + { { "%8.3f" }, "10.3456" }, - //(Minimum)Eight-wide,two positions after the decimal,zero-filled,default(right)-justified + //(Minimum)Eight-wide,two positions after the + // decimal,zero-filled,default(right)-justified - {"%08.2f","10.3456"}, + { { "%08.2f" }, "10.3456" }, //(Minimum)Eight-wide,two positions after the decimal,left-justified - {"%-8.2f","10.3456"}, + { { "%-8.2f" }, "10.3456" }, - //(Minimum)Eight-wide,two positions after the decimal,with sign,default(right)-justified + //(Minimum)Eight-wide,two positions after the decimal,with + // sign,default(right)-justified - {"%+8.2f","-10.3456"}, + { { "%+8.2f" }, "-10.3456" }, - //Zero positions after the decimal([floor]rounding),default(right)-justified + // Zero positions after the + // decimal([floor]rounding),default(right)-justified - {"%.0f","0.1"}, + { { "%.0f" }, "0.1" }, - //Zero positions after the decimal([ceil]rounding),default(right)-justified + // Zero positions after the decimal([ceil]rounding),default(right)-justified - {"%.0f","0.6"}, + { { "%.0f" }, "0.6" }, - //Zero-filled,default positions number after the decimal,default(right)-justified + // Zero-filled,default positions number after the + // decimal,default(right)-justified - {"%0f","0.6"}, + { { "%0f" }, "0.6" }, - //Double argument representing floating-point,used by f style,default(right)-justified + // Double argument representing floating-point,used by f + // style,default(right)-justified - {"%4g","12345.6789"}, + { { "%4g" }, "12345.6789" }, - //Double argument representing floating-point,used by e style,default(right)-justified + // Double argument representing floating-point,used by e + // style,default(right)-justified - {"%4.2g","12345.6789"}, + { { "%4.2g" }, "12345.6789" }, - //Double argument representing floating-point,used by f style,default(right)-justified + // Double argument representing floating-point,used by f + // style,default(right)-justified - {"%4G","0.0000023"}, + { { "%4G" }, "0.0000023" }, - //Double argument representing floating-point,used by e style,default(right)-justified + // Double argument representing floating-point,used by e + // style,default(right)-justified - {"%4G","0.023"}, + { { "%4G" }, "0.023" }, - //Double argument representing floating-point,with exponent,left-justified,default(right)-justified + // Double argument representing floating-point,with + // exponent,left-justified,default(right)-justified - {"%-#20.15e","789456123.0"}, + { { "%-#20.15e" }, "789456123.0" }, - //Double argument representing floating-point,with exponent,left-justified,with sign,capital E,default(right)-justified + // Double argument representing floating-point,with + // exponent,left-justified,with sign,capital E,default(right)-justified - {"%+#21.15E","789456123.0"}, + { { "%+#21.15E" }, "789456123.0" }, - //Double argument representing floating-point,in [-]xh.hhhhpAd style + // Double argument representing floating-point,in [-]xh.hhhhpAd style - {"%.6a","0.1"}, + { { "%.6a" }, "0.1" }, - //(Minimum)Ten-wide,Double argument representing floating-point,in xh.hhhhpAd style,default(right)-justified + //(Minimum)Ten-wide,Double argument representing floating-point,in + // xh.hhhhpAd style,default(right)-justified - {"%10.2a","9990.235"}, + { { "%10.2a" }, "9990.235" }, }; //--------------------------------------------------------- @@ -377,14 +391,14 @@ std::vector printFloatLimitsGenParameters = { // Infinity (1.0/0.0) - { "%f", "1.0f/0.0f" }, + { { "%f" }, "1.0f/0.0f" }, // NaN - { "%f", "sqrt(-1.0f)" }, + { { "%f" }, "sqrt(-1.0f)" }, // NaN - { "%f", "acospi(2.0f)" } + { { "%f" }, "acospi(2.0f)" } }; //-------------------------------------------------------- @@ -437,21 +451,22 @@ testCase testCaseFloatLimits = { std::vector printOctalGenParameters = { - //Default(right)-justified + // Default(right)-justified - {"%o","10"}, + { { "%o" }, "10" }, - //Five-digit,default(right)-justified + // Five-digit,default(right)-justified - {"%.5o","10"}, + { { "%.5o" }, "10" }, - //Default(right)-justified,increase precision + // Default(right)-justified,increase precision - {"%#o","100000000"}, + { { "%#o" }, "100000000" }, - //(Minimum)Four-wide,Five-digit,0-flag ignored(because of precision),default(right)-justified + //(Minimum)Four-wide,Five-digit,0-flag ignored(because of + // precision),default(right)-justified - {"%04.5o","10"} + { { "%04.5o" }, "10" } }; @@ -493,19 +508,19 @@ testCase testCaseOctal = { std::vector printUnsignedGenParameters = { - //Default(right)-justified + // Default(right)-justified - {"%u","10"}, + { { "%u" }, "10" }, - //Zero precision for zero,default(right)-justified + // Zero precision for zero,default(right)-justified - {"%.0u","0"}, + { { "%.0u" }, "0" }, }; //------------------------------------------------------- -//Test case for octal | +// Test case for unsigned | //------------------------------------------------------- @@ -541,25 +556,25 @@ testCase testCaseUnsigned = { std::vector printHexadecimalGenParameters = { - //Add 0x,low x,default(right)-justified + // Add 0x,low x,default(right)-justified - {"%#x","0xABCDEF"}, + { { "%#x" }, "0xABCDEF" }, - //Add 0x,capital X,default(right)-justified + // Add 0x,capital X,default(right)-justified - {"%#X","0xABCDEF"}, + { { "%#X" }, "0xABCDEF" }, - //Not add 0x,if zero,default(right)-justified + // Not add 0x,if zero,default(right)-justified - {"%#X","0"}, + { { "%#X" }, "0" }, //(Minimum)Eight-wide,default(right)-justified - {"%8x","399"}, + { { "%8x" }, "399" }, //(Minimum)Four-wide,zero-filled,default(right)-justified - {"%04x","399"} + { { "%04x" }, "399" } }; @@ -601,17 +616,17 @@ testCase testCaseHexadecimal = { std::vector printCharGenParameters = { - //Four-wide,zero-filled,default(right)-justified + // Four-wide,zero-filled,default(right)-justified - {"%4c","\'1\'"}, + { { "%4c" }, "\'1\'" }, - //Four-wide,left-justified + // Four-wide,left-justified - {"%-4c","\'1\'"}, + { { "%-4c" }, "\'1\'" }, - //(unsigned) int argument,default(right)-justified + //(unsigned) int argument,default(right)-justified - {"%c","66"} + { { "%c" }, "66" } }; @@ -669,20 +684,49 @@ testCase testCaseChar = { // [string]format | [string] string-data representation | //-------------------------------------------------------- +// clang-format off std::vector printStringGenParameters = { //(Minimum)Four-wide,zero-filled,default(right)-justified - {"%4s","\"foo\""}, + { { "%4s" }, "\"foo\"" }, - //One-digit(precision ignored),left-justified + // One-digit(precision ignored),left-justified - {"%.1s","\"foo\""}, + { { "%.1s" }, "\"foo\"" }, //%% specification - {"%s","\"%%\""}, + { { "%s" }, "\"%%\"" }, + + // special symbols + // nested + + { { "%s" }, "\"\\\"%%\\\"\"" }, + + { { "%s" }, "\"\\\'%%\\\'\"" }, + + // tabs + + { { "%s" }, "\"foo\\tfoo\"" }, + + // newlines + + { { "%s" }, "\"foo\\nfoo\"" }, + + // terminator + { { "%s" }, "\"foo\\0foo\"" }, + + // carriage return + { { "%s" }, "\"foo\\rfoo\"" }, + + // all ascii characters + { { "%s" }, + "\" " + "!\\\"#$%&\'()*+,-./" + "0123456789:;<=>?@ABCDEFGHIJKLMNOPQRSTUVWXYZ[\\\\]^_`" + "abcdefghijklmnopqrstuvwxyz{|}~\"" } }; //--------------------------------------------------------- @@ -698,8 +742,24 @@ std::vector correctBufferString = { "f", "%%", -}; + "\"%%\"", + + "\'%%\'", + + "foo\tfoo", + +R"(foo +foo)", + + "foo", + + "foo", + + " !\"#$%&\'()*+,-./" + "0123456789:;<=>?@ABCDEFGHIJKLMNOPQRSTUVWXYZ[\\]^_`" + "abcdefghijklmnopqrstuvwxyz{|}~" +}; //--------------------------------------------------------- @@ -721,7 +781,84 @@ testCase testCaseString = { }; +//-------------------------------------------------------- + +// [string]format | + +//-------------------------------------------------------- + +std::vector printFormatStringGenParameters = { + + //%% specification + + { { "%%" } }, + + // special symbols + // nested + + { { "\\\"%%\\\"" } }, + + { { "\'%%\'" } }, + + // tabs + + { { "foo\\t\\t\\tfoo" } }, + + // newlines + + { { "foo\\nfoo" } }, + + // all ascii characters + { { " !\\\"#$%&\'()*+,-./" + "0123456789:;<=>?@ABCDEFGHIJKLMNOPQRSTUVWXYZ[\\\\]^_`" + "abcdefghijklmnopqrstuvwxyz{|}~" } } +}; + +//--------------------------------------------------------- +// Lookup table -[string] string-correct buffer | + +//--------------------------------------------------------- + +std::vector correctBufferFormatString = { + + "%", + + "\"%\"", + + "\'%\'", + + "foo\t\t\tfoo", + +R"(foo +foo)", + + " !\"#$%&\'()*+,-./" + "0123456789:;<=>?@ABCDEFGHIJKLMNOPQRSTUVWXYZ[\\]^_`" + "abcdefghijklmnopqrstuvwxyz{|}~" +}; + +//--------------------------------------------------------- + +//Test case for string | + +//--------------------------------------------------------- + +testCase testCaseFormatString = { + + TYPE_FORMAT_STRING, + + correctBufferFormatString, + + printFormatStringGenParameters, + + NULL, + + kchar + +}; + +// clang-format on //========================================================= @@ -741,27 +878,27 @@ std::vector printVectorGenParameters = { //(Minimum)Two-wide,two positions after decimal - { NULL, "(1.0f,2.0f,3.0f,4.0f)", "%2.2", "hlf", "float", "4" }, + { {}, "(1.0f,2.0f,3.0f,4.0f)", "%2.2", "hlf", "float", "4" }, // Alternative form,uchar argument - { NULL, "(0xFA,0xFB)", "%#", "hhx", "uchar", "2" }, + { {}, "(0xFA,0xFB)", "%#", "hhx", "uchar", "2" }, // Alternative form,ushort argument - { NULL, "(0x1234,0x8765)", "%#", "hx", "ushort", "2" }, + { {}, "(0x1234,0x8765)", "%#", "hx", "ushort", "2" }, // Alternative form,uint argument - { NULL, "(0x12345678,0x87654321)", "%#", "hlx", "uint", "2" }, + { {}, "(0x12345678,0x87654321)", "%#", "hlx", "uint", "2" }, // Alternative form,long argument - { NULL, "(12345678,98765432)", "%", "ld", "long", "2" }, + { {}, "(12345678,98765432)", "%", "ld", "long", "2" }, //(Minimum)Two-wide,two positions after decimal - { NULL, "(1.0h,2.0h,3.0h,4.0h)", "%2.2", "hf", "half", "4" } + { {}, "(1.0h,2.0h,3.0h,4.0h)", "%2.2", "hf", "half", "4" } }; //------------------------------------------------------------ @@ -821,28 +958,72 @@ testCase testCaseVector = { //------------------------------------------------------------------------------------------------------------------------------------------------------------------- - std::vector printAddrSpaceGenParameters = { - //Global memory region - - {"\"%d\\n\"",NULL,NULL,NULL,NULL,NULL,"__global int* x","","*x",""}, - - //Global,constant, memory region - - {"\"%d\\n\"",NULL,NULL,NULL,NULL,NULL,"constant int* x","","*x",""}, - - //Local memory region - - {"\"%+d\\n\"",NULL,NULL,NULL,NULL,NULL,"","local int x;\n x= (int)3;\n","x",""}, - - //Private memory region - - {"\"%i\\n\"",NULL,NULL,NULL,NULL,NULL,"","private int x;\n x = (int)-1;\n","x",""}, - - //Address of void * from global memory region - - {"\"%p\\n\"",NULL,NULL,NULL,NULL,NULL,"__global void* x,__global intptr_t* xAddr","","x","*xAddr = (intptr_t)x;\n"} + // Global memory region + + { { "\"%d\\n\"" }, + NULL, + NULL, + NULL, + NULL, + NULL, + "__global int* x", + "", + "*x", + "" }, + + // Global,constant, memory region + + { { "\"%d\\n\"" }, + NULL, + NULL, + NULL, + NULL, + NULL, + "constant int* x", + "", + "*x", + "" }, + + // Local memory region + + { { "\"%+d\\n\"" }, + NULL, + NULL, + NULL, + NULL, + NULL, + "", + "local int x;\n x= (int)3;\n", + "x", + "" }, + + // Private memory region + + { { "\"%i\\n\"" }, + NULL, + NULL, + NULL, + NULL, + NULL, + "", + "private int x;\n x = (int)-1;\n", + "x", + "" }, + + // Address of void * from global memory region + + { { "\"%p\\n\"" }, + NULL, + NULL, + NULL, + NULL, + NULL, + "__global void* x,__global intptr_t* xAddr", + "", + "x", + "*xAddr = (intptr_t)x;\n" } }; @@ -885,10 +1066,11 @@ testCase testCaseAddrSpace = { //------------------------------------------------------------------------------- std::vector allTestCase = { - &testCaseInt, &testCaseHalf, &testCaseHalfLimits, - &testCaseFloat, &testCaseFloatLimits, &testCaseOctal, - &testCaseUnsigned, &testCaseHexadecimal, &testCaseChar, - &testCaseString, &testCaseVector, &testCaseAddrSpace + &testCaseInt, &testCaseHalf, &testCaseHalfLimits, + &testCaseFloat, &testCaseFloatLimits, &testCaseOctal, + &testCaseUnsigned, &testCaseHexadecimal, &testCaseChar, + &testCaseString, &testCaseFormatString, &testCaseVector, + &testCaseAddrSpace }; //----------------------------------------- @@ -920,7 +1102,7 @@ size_t verifyOutputBuffer(char *analysisBuffer,testCase* pTestCase,size_t testId strcpy(analysisBufferTmp,"0x"); else analysisBufferTmp[0]='\0'; - strcat(analysisBufferTmp,analysisBuffer); + strncat(analysisBufferTmp, analysisBuffer, ANALYSIS_BUFFER_SIZE); if (sizeof(long) == 8) { if(strtoul(analysisBufferTmp,NULL,0) == pAddr) return 0; } @@ -965,7 +1147,8 @@ size_t verifyOutputBuffer(char *analysisBuffer,testCase* pTestCase,size_t testId static void intRefBuilder(printDataGenParameters& params, char* refResult, const size_t refSize) { - snprintf(refResult, refSize, params.genericFormat, atoi(params.dataRepresentation)); + snprintf(refResult, refSize, params.genericFormats.front().c_str(), + atoi(params.dataRepresentation)); } static void halfRefBuilder(printDataGenParameters& params, char* refResult, @@ -973,30 +1156,32 @@ static void halfRefBuilder(printDataGenParameters& params, char* refResult, { cl_half val = cl_half_from_float(strtof(params.dataRepresentation, NULL), half_rounding_mode); - snprintf(refResult, refSize, params.genericFormat, cl_half_to_float(val)); + snprintf(refResult, refSize, params.genericFormats.front().c_str(), + cl_half_to_float(val)); } static void floatRefBuilder(printDataGenParameters& params, char* refResult, const size_t refSize) { - snprintf(refResult, refSize, params.genericFormat, strtof(params.dataRepresentation, NULL)); + snprintf(refResult, refSize, params.genericFormats.front().c_str(), + strtof(params.dataRepresentation, NULL)); } static void octalRefBuilder(printDataGenParameters& params, char* refResult, const size_t refSize) { const unsigned long int data = strtoul(params.dataRepresentation, NULL, 10); - snprintf(refResult, refSize, params.genericFormat, data); + snprintf(refResult, refSize, params.genericFormats.front().c_str(), data); } static void unsignedRefBuilder(printDataGenParameters& params, char* refResult, const size_t refSize) { const unsigned long int data = strtoul(params.dataRepresentation, NULL, 10); - snprintf(refResult, refSize, params.genericFormat, data); + snprintf(refResult, refSize, params.genericFormats.front().c_str(), data); } static void hexRefBuilder(printDataGenParameters& params, char* refResult, const size_t refSize) { const unsigned long int data = strtoul(params.dataRepresentation, NULL, 0); - snprintf(refResult, refSize, params.genericFormat, data); + snprintf(refResult, refSize, params.genericFormats.front().c_str(), data); } /* From 61ae64301743b8aa3f313d20c3100c98f3856347 Mon Sep 17 00:00:00 2001 From: Marcin Hajder Date: Fri, 19 Apr 2024 14:03:49 +0200 Subject: [PATCH 2/7] removed one case due to compatibility concerns --- test_conformance/printf/util_printf.cpp | 5 ----- 1 file changed, 5 deletions(-) diff --git a/test_conformance/printf/util_printf.cpp b/test_conformance/printf/util_printf.cpp index deb627958..e176ecc9c 100644 --- a/test_conformance/printf/util_printf.cpp +++ b/test_conformance/printf/util_printf.cpp @@ -718,9 +718,6 @@ std::vector printStringGenParameters = { // terminator { { "%s" }, "\"foo\\0foo\"" }, - // carriage return - { { "%s" }, "\"foo\\rfoo\"" }, - // all ascii characters { { "%s" }, "\" " @@ -754,8 +751,6 @@ foo)", "foo", - "foo", - " !\"#$%&\'()*+,-./" "0123456789:;<=>?@ABCDEFGHIJKLMNOPQRSTUVWXYZ[\\]^_`" "abcdefghijklmnopqrstuvwxyz{|}~" From 40d5a835b9e1fbe8eb7af26338bd58558ce6230f Mon Sep 17 00:00:00 2001 From: Marcin Hajder Date: Fri, 19 Apr 2024 14:32:22 +0200 Subject: [PATCH 3/7] warning correction --- test_conformance/printf/util_printf.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/test_conformance/printf/util_printf.cpp b/test_conformance/printf/util_printf.cpp index e176ecc9c..07f4733b4 100644 --- a/test_conformance/printf/util_printf.cpp +++ b/test_conformance/printf/util_printf.cpp @@ -1090,7 +1090,7 @@ size_t verifyOutputBuffer(char *analysisBuffer,testCase* pTestCase,size_t testId if(pTestCase->_type == TYPE_ADDRESS_SPACE && strcmp(pTestCase->_genParameters[testId].addrSpacePAdd,"")) { - char analysisBufferTmp[ANALYSIS_BUFFER_SIZE]; + char analysisBufferTmp[ANALYSIS_BUFFER_SIZE+1]; if(strstr(analysisBuffer,"0x") == NULL) // Need to prepend 0x to ASCII number before calling strtol. From bd4ec5eba01b42f5c02a888ceb2d46f6e1ddd036 Mon Sep 17 00:00:00 2001 From: Marcin Hajder Date: Fri, 19 Apr 2024 15:33:55 +0200 Subject: [PATCH 4/7] clang format --- test_conformance/printf/util_printf.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/test_conformance/printf/util_printf.cpp b/test_conformance/printf/util_printf.cpp index 07f4733b4..3bcca0267 100644 --- a/test_conformance/printf/util_printf.cpp +++ b/test_conformance/printf/util_printf.cpp @@ -1090,7 +1090,7 @@ size_t verifyOutputBuffer(char *analysisBuffer,testCase* pTestCase,size_t testId if(pTestCase->_type == TYPE_ADDRESS_SPACE && strcmp(pTestCase->_genParameters[testId].addrSpacePAdd,"")) { - char analysisBufferTmp[ANALYSIS_BUFFER_SIZE+1]; + char analysisBufferTmp[ANALYSIS_BUFFER_SIZE + 1]; if(strstr(analysisBuffer,"0x") == NULL) // Need to prepend 0x to ASCII number before calling strtol. From b5cef1c934aab26fca1698dbe8cf32fbb5da4135 Mon Sep 17 00:00:00 2001 From: Marcin Hajder Date: Fri, 19 Apr 2024 17:48:59 +0200 Subject: [PATCH 5/7] fixed warning --- test_conformance/printf/test_printf.cpp | 11 ++++++----- 1 file changed, 6 insertions(+), 5 deletions(-) diff --git a/test_conformance/printf/test_printf.cpp b/test_conformance/printf/test_printf.cpp index 0c5eb169a..1dea536ed 100644 --- a/test_conformance/printf/test_printf.cpp +++ b/test_conformance/printf/test_printf.cpp @@ -162,12 +162,13 @@ static void getAnalysisBuffer(char* analysisBuffer) FILE *fp; memset(analysisBuffer,0,ANALYSIS_BUFFER_SIZE); - fp = fopen(gFileName,"r"); - if(NULL == fp) + fp = fopen(gFileName, "r"); + if (NULL == fp) log_error("Failed to open analysis buffer ('%s')\n", strerror(errno)); - else - std::fread(analysisBuffer, sizeof(analysisBuffer[0]), - ANALYSIS_BUFFER_SIZE, fp); + else if (0 + == std::fread(analysisBuffer, sizeof(analysisBuffer[0]), + ANALYSIS_BUFFER_SIZE, fp)) + log_error("Failed to read analysis buffer\n"); fclose(fp); } From c39a288e1ac25e59e884a95494091fe0ecd1036f Mon Sep 17 00:00:00 2001 From: Marcin Hajder Date: Thu, 16 May 2024 21:39:45 +0200 Subject: [PATCH 6/7] Corrections related to code review process --- test_conformance/printf/test_printf.cpp | 430 +++++++++++------------- test_conformance/printf/test_printf.h | 2 +- test_conformance/printf/util_printf.cpp | 228 ++++++------- 3 files changed, 298 insertions(+), 362 deletions(-) diff --git a/test_conformance/printf/test_printf.cpp b/test_conformance/printf/test_printf.cpp index 1dea536ed..02d6660f1 100644 --- a/test_conformance/printf/test_printf.cpp +++ b/test_conformance/printf/test_printf.cpp @@ -16,6 +16,7 @@ #include "harness/os_helpers.h" #include "harness/typeWrappers.h" +#include #include #include #include @@ -82,8 +83,7 @@ static int isKernelPFormat(testCase* pTestCase,size_t testId); 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); + bool isLongSupport = true, bool is64bAddrSpace = false); // Creates and execute the printf test for the given device, context, type/format static int doTest(cl_command_queue queue, cl_context context, @@ -168,7 +168,7 @@ static void getAnalysisBuffer(char* analysisBuffer) else if (0 == std::fread(analysisBuffer, sizeof(analysisBuffer[0]), ANALYSIS_BUFFER_SIZE, fp)) - log_error("Failed to read analysis buffer\n"); + log_error("No data read from analysis buffer\n"); fclose(fp); } @@ -220,7 +220,6 @@ 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) { int err; @@ -265,10 +264,7 @@ static cl_program makePrintfProgram(cl_kernel* kernel_ptr, ->_genParameters[testNum] .addrSpaceVariableTypeQualifier, "printf(", - allTestCase[testId] - ->_genParameters[testNum] - .genericFormats[formatNum] - .c_str(), + allTestCase[testId]->_genParameters[testNum].genericFormat, ",", allTestCase[testId]->_genParameters[testNum].addrSpaceParameter, "); ", @@ -329,10 +325,7 @@ static cl_program makePrintfProgram(cl_kernel* kernel_ptr, << "(void)\n" "{\n" " printf(\"" - << allTestCase[testId] - ->_genParameters[testNum] - .genericFormats[formatNum] - .c_str() + << allTestCase[testId]->_genParameters[testNum].genericFormat << "\\n\""; if (allTestCase[testId]->_genParameters[testNum].dataRepresentation) @@ -343,7 +336,7 @@ static cl_program makePrintfProgram(cl_kernel* kernel_ptr, .dataRepresentation; } - sourceGen << ");}\n"; + sourceGen << ");\n}\n"; std::string kernel_source = sourceGen.str(); const char* ptr = kernel_source.c_str(); @@ -433,12 +426,18 @@ static bool is64bAddressSpace(cl_device_id device_id) //----------------------------------------- // subtest_fail //----------------------------------------- -template void subtest_fail(const char* msg, Args... args) +void subtest_fail(const char* msg, ...) { - if (msg) log_error(msg, args...); + if (msg) + { + va_list argptr; + va_start(argptr, msg); + vfprintf(stderr, msg, argptr); + va_end(argptr); + } ++s_test_fail; ++s_test_cnt; -}; +} //----------------------------------------- // doTest @@ -465,257 +464,234 @@ static int doTest(cl_command_queue queue, cl_context context, for (unsigned testNum = 0; testNum < genParams.size(); testNum++) { - for (unsigned formatNum = 0; formatNum < allTestCase[testId] - ->_genParameters[testNum] - .genericFormats.size(); - formatNum++) + if (allTestCase[testId]->_type == TYPE_VECTOR) { - if (allTestCase[testId]->_type == TYPE_VECTOR) + if ((strcmp(allTestCase[testId]->_genParameters[testNum].dataType, + "half") + == 0) + && !is_extension_available(device, "cl_khr_fp16")) { - if ((strcmp( - allTestCase[testId]->_genParameters[testNum].dataType, - "half") - == 0) - && !is_extension_available(device, "cl_khr_fp16")) - { - log_info( - "Skipping half because cl_khr_fp16 extension is not " - "supported.\n"); + log_info("Skipping half because cl_khr_fp16 extension is not " + "supported.\n"); - s_test_skip++; - s_test_cnt++; - continue; - } + s_test_skip++; + s_test_cnt++; + continue; + } + 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 printf(\"%sv%s%s\",%s)\n", testNum, + "%d)testing kernel //argument %s \n printf(%s,%s)\n", + testNum, allTestCase[testId] ->_genParameters[testNum] - .vectorFormatFlag, - allTestCase[testId]->_genParameters[testNum].vectorSize, + .addrSpaceArgumentTypeQualifier, + allTestCase[testId]->_genParameters[testNum].genericFormat, 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); - } + .addrSpaceParameter); } else { - log_info("%d)testing printf(\"%s\",%s)\n", testNum, - allTestCase[testId] - ->_genParameters[testNum] - .genericFormats[formatNum] - .c_str(), - allTestCase[testId] - ->_genParameters[testNum] - .dataRepresentation); + log_info( + "%d)testing kernel //variable %s \n printf(%s,%s)\n", + testNum, + allTestCase[testId] + ->_genParameters[testNum] + .addrSpaceVariableTypeQualifier, + allTestCase[testId]->_genParameters[testNum].genericFormat, + allTestCase[testId] + ->_genParameters[testNum] + .addrSpaceParameter); } + } + else + { + log_info("%d)testing printf(\"%s\",%s)\n", testNum, + allTestCase[testId]->_genParameters[testNum].genericFormat, + allTestCase[testId] + ->_genParameters[testNum] + .dataRepresentation); + } - // Long support for varible type - if (allTestCase[testId]->_type == TYPE_VECTOR - && !strcmp( - allTestCase[testId]->_genParameters[testNum].dataType, - "long") - && !isLongSupported(device)) - { - log_info("Long is not supported, test not run.\n"); - s_test_skip++; - s_test_cnt++; - continue; - } + // Long support for varible type + if (allTestCase[testId]->_type == TYPE_VECTOR + && !strcmp(allTestCase[testId]->_genParameters[testNum].dataType, + "long") + && !isLongSupported(device)) + { + log_info("Long is not supported, test not run.\n"); + s_test_skip++; + s_test_cnt++; + continue; + } - fflush(stdout); + 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; - } + // 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; + } - clProgramWrapper program; - clKernelWrapper kernel; - clMemWrapper d_out; - clMemWrapper d_a; - char _analysisBuffer[ANALYSIS_BUFFER_SIZE]; - cl_uint out32 = 0; - cl_ulong out64 = 0; - int fd = -1; - - // Define an index space (global work size) of threads for - // execution. - size_t globalWorkSize[1]; - - program = - makePrintfProgram(&kernel, context, testId, testNum, formatNum, - isLongSupport, is64bAddressSpace(device)); - if (!program || !kernel) - { - subtest_fail(nullptr); - continue; - } + clProgramWrapper program; + clKernelWrapper kernel; + clMemWrapper d_out; + clMemWrapper d_a; + char _analysisBuffer[ANALYSIS_BUFFER_SIZE]; + cl_uint out32 = 0; + cl_ulong out64 = 0; + int fd = -1; + + // Define an index space (global work size) of threads for + // execution. + size_t globalWorkSize[1]; + + program = makePrintfProgram(&kernel, context, testId, testNum, + isLongSupport, is64bAddressSpace(device)); + if (!program || !kernel) + { + subtest_fail(nullptr); + continue; + } - // For address space test if there is kernel argument - set it - if (allTestCase[testId]->_type == TYPE_ADDRESS_SPACE) + // For address space test if there is kernel argument - set it + if (allTestCase[testId]->_type == TYPE_ADDRESS_SPACE) + { + if (isKernelArgument(allTestCase[testId], testNum)) { - if (isKernelArgument(allTestCase[testId], testNum)) + int a = 2; + d_a = clCreateBuffer(context, + CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, + sizeof(int), &a, &err); + if (err != CL_SUCCESS || d_a == NULL) { - int a = 2; - d_a = clCreateBuffer( - context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, - sizeof(int), &a, &err); - if (err != CL_SUCCESS || d_a == NULL) - { - subtest_fail("clCreateBuffer failed\n"); - continue; - } - err = clSetKernelArg(kernel, 0, sizeof(cl_mem), &d_a); - if (err != CL_SUCCESS) - { - subtest_fail("clSetKernelArg failed\n"); - continue; - } + subtest_fail("clCreateBuffer failed\n"); + continue; } - // For address space test if %p is tested - if (isKernelPFormat(allTestCase[testId], testNum)) + err = clSetKernelArg(kernel, 0, sizeof(cl_mem), &d_a); + if (err != CL_SUCCESS) { - d_out = clCreateBuffer(context, CL_MEM_READ_WRITE, - sizeof(cl_ulong), NULL, &err); - if (err != CL_SUCCESS || d_out == NULL) - { - subtest_fail("clCreateBuffer failed\n"); - continue; - } - err = clSetKernelArg(kernel, 1, sizeof(cl_mem), &d_out); - if (err != CL_SUCCESS) - { - subtest_fail("clSetKernelArg failed\n"); - continue; - } + subtest_fail("clSetKernelArg failed\n"); + continue; } } - - fd = acquireOutputStream(&err); - if (err != 0) - { - subtest_fail("Error while redirection stdout to file"); - continue; - } - globalWorkSize[0] = 1; - cl_event ndrEvt; - err = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, globalWorkSize, - NULL, 0, NULL, &ndrEvt); - if (err != CL_SUCCESS) + // For address space test if %p is tested + if (isKernelPFormat(allTestCase[testId], testNum)) { - releaseOutputStream(fd); - subtest_fail("\n clEnqueueNDRangeKernel failed errcode:%d\n", - err); - continue; + d_out = clCreateBuffer(context, CL_MEM_READ_WRITE, + sizeof(cl_ulong), NULL, &err); + if (err != CL_SUCCESS || d_out == NULL) + { + subtest_fail("clCreateBuffer failed\n"); + continue; + } + err = clSetKernelArg(kernel, 1, sizeof(cl_mem), &d_out); + if (err != CL_SUCCESS) + { + subtest_fail("clSetKernelArg failed\n"); + continue; + } } + } - fflush(stdout); - err = clFlush(queue); - if (err != CL_SUCCESS) - { - releaseOutputStream(fd); - subtest_fail("clFlush failed : %d\n", err); - continue; - } - // Wait until kernel finishes its execution and (thus) the output - // printed from the kernel is immediately printed - err = waitForEvent(&ndrEvt); + fd = acquireOutputStream(&err); + if (err != 0) + { + subtest_fail("Error while redirection stdout to file"); + continue; + } + globalWorkSize[0] = 1; + cl_event ndrEvt; + err = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, globalWorkSize, + NULL, 0, NULL, &ndrEvt); + if (err != CL_SUCCESS) + { + releaseOutputStream(fd); + subtest_fail("\n clEnqueueNDRangeKernel failed errcode:%d\n", err); + continue; + } + fflush(stdout); + err = clFlush(queue); + if (err != CL_SUCCESS) + { releaseOutputStream(fd); + subtest_fail("clFlush failed : %d\n", err); + continue; + } + // Wait until kernel finishes its execution and (thus) the output + // printed from the kernel is immediately printed + err = waitForEvent(&ndrEvt); + + releaseOutputStream(fd); - if (err != CL_SUCCESS) + if (err != CL_SUCCESS) + { + subtest_fail("waitforEvent failed : %d\n", err); + continue; + } + fflush(stdout); + + if (allTestCase[testId]->_type == TYPE_ADDRESS_SPACE + && isKernelPFormat(allTestCase[testId], testNum)) + { + // Read the OpenCL output buffer (d_out) to the host output + // array (out) + if (!is64bAddressSpace(device)) // 32-bit address space { - subtest_fail("waitforEvent failed : %d\n", err); - continue; + clEnqueueReadBuffer(queue, d_out, CL_TRUE, 0, sizeof(cl_int), + &out32, 0, NULL, NULL); } - fflush(stdout); - - if (allTestCase[testId]->_type == TYPE_ADDRESS_SPACE - && isKernelPFormat(allTestCase[testId], testNum)) + else // 64-bit address space { - // Read the OpenCL output buffer (d_out) to the host output - // array (out) - if (!is64bAddressSpace(device)) // 32-bit address space - { - clEnqueueReadBuffer(queue, d_out, CL_TRUE, 0, - sizeof(cl_int), &out32, 0, NULL, NULL); - } - else // 64-bit address space - { - clEnqueueReadBuffer(queue, d_out, CL_TRUE, 0, - sizeof(cl_ulong), &out64, 0, NULL, - NULL); - } + clEnqueueReadBuffer(queue, d_out, CL_TRUE, 0, sizeof(cl_ulong), + &out64, 0, NULL, NULL); } + } - // - // Get the output printed from the kernel to _analysisBuffer - // and verify its correctness - getAnalysisBuffer(_analysisBuffer); - if (!is64bAddressSpace(device)) // 32-bit address space + // + // Get the output printed from the kernel to _analysisBuffer + // and verify its correctness + getAnalysisBuffer(_analysisBuffer); + if (!is64bAddressSpace(device)) // 32-bit address space + { + if (0 + != verifyOutputBuffer(_analysisBuffer, allTestCase[testId], + testNum, (cl_ulong)out32)) { - if (0 - != verifyOutputBuffer(_analysisBuffer, allTestCase[testId], - testNum, (cl_ulong)out32)) - { - subtest_fail("verifyOutputBuffer failed\n"); - continue; - } + subtest_fail("verifyOutputBuffer failed\n"); + continue; } - else // 64-bit address space + } + else // 64-bit address space + { + if (0 + != verifyOutputBuffer(_analysisBuffer, allTestCase[testId], + testNum, out64)) { - if (0 - != verifyOutputBuffer(_analysisBuffer, allTestCase[testId], - testNum, out64)) - { - subtest_fail("verifyOutputBuffer failed\n"); - continue; - } + subtest_fail("verifyOutputBuffer failed\n"); + continue; } - ++s_test_cnt; } + ++s_test_cnt; } // all subtests skipped ? diff --git a/test_conformance/printf/test_printf.h b/test_conformance/printf/test_printf.h index 0a33d5f84..0eca62bdb 100644 --- a/test_conformance/printf/test_printf.h +++ b/test_conformance/printf/test_printf.h @@ -63,7 +63,7 @@ enum PrintfTestType struct printDataGenParameters { - std::vector genericFormats; + const char* genericFormat; const char* dataRepresentation; const char* vectorFormatFlag; const char* vectorFormatSpecifier; diff --git a/test_conformance/printf/util_printf.cpp b/test_conformance/printf/util_printf.cpp index 3bcca0267..67f11d09e 100644 --- a/test_conformance/printf/util_printf.cpp +++ b/test_conformance/printf/util_printf.cpp @@ -46,41 +46,41 @@ std::vector printIntGenParameters = { //(Minimum)Five-wide,default(right)-justified - { { "%5d" }, "10" }, + { "%5d", "10" }, //(Minimum)Five-wide,left-justified - { { "%-5d" }, "10" }, + { "%-5d", "10" }, //(Minimum)Five-wide,default(right)-justified,zero-filled - { { "%05d" }, "10" }, + { "%05d", "10" }, //(Minimum)Five-wide,default(right)-justified,with sign - { { "%+5d" }, "10" }, + { "%+5d", "10" }, //(Minimum)Five-wide ,left-justified,with sign - { { "%-+5d" }, "10" }, + { "%-+5d", "10" }, //(Minimum)Five-digit(zero-filled in absent digits),default(right)-justified - { { "%.5i" }, "100" }, + { "%.5i", "100" }, //(Minimum)Six-wide,Five-digit(zero-filled in absent // digits),default(right)-justified - { { "%6.5i" }, "100" }, + { "%6.5i", "100" }, // 0 and - flag both apper ==>0 is ignored,left-justified,capital I - { { "%-06i" }, "100" }, + { "%-06i", "100" }, //(Minimum)Six-wide,Five-digit(zero-filled in absent // digits),default(right)-justified - { { "%06.5i" }, "100" } + { "%06.5i", "100" } }; @@ -121,50 +121,50 @@ std::vector printHalfGenParameters = { // Default(right)-justified - { { "%f" }, "1.234h" }, + { "%f", "1.234h" }, // One position after the decimal,default(right)-justified - { { "%4.2f" }, "1.2345h" }, + { "%4.2f", "1.2345h" }, // Zero positions after the // decimal([floor]rounding),default(right)-justified - { { "%.0f" }, "0.1h" }, + { "%.0f", "0.1h" }, // Zero positions after the decimal([ceil]rounding),default(right)-justified - { { "%.0f" }, "0.6h" }, + { "%.0f", "0.6h" }, // Zero-filled,default positions number after the // decimal,default(right)-justified - { { "%0f" }, "0.6h" }, + { "%0f", "0.6h" }, // Double argument representing floating-point,used by f // style,default(right)-justified - { { "%4g" }, "5.678h" }, + { "%4g", "5.678h" }, // Double argument representing floating-point,used by e // style,default(right)-justified - { { "%4.2g" }, "5.678h" }, + { "%4.2g", "5.678h" }, // Double argument representing floating-point,used by e // style,default(right)-justified - { { "%4G" }, "0.000062h" }, + { "%4G", "0.000062h" }, // Double argument representing floating-point,with // exponent,left-justified,default(right)-justified - { { "%-#20.15e" }, "65504.0h" }, + { "%-#20.15e", "65504.0h" }, // Double argument representing floating-point,with // exponent,left-justified,with sign,capital E,default(right)-justified - { { "%+#21.15E" }, "-65504.0h" }, + { "%+#21.15E", "-65504.0h" }, }; //--------------------------------------------------------- @@ -206,14 +206,14 @@ std::vector printHalfLimitsGenParameters = { // Infinity (1.0/0.0) - { { "%f" }, "1.0h/0.0h" }, + { "%f", "1.0h/0.0h" }, // NaN - { { "%f" }, "sqrt(-1.0h)" }, + { "%f", "sqrt(-1.0h)" }, // NaN - { { "%f" }, "acospi(2.0h)" } + { "%f", "acospi(2.0h)" } }; //-------------------------------------------------------- @@ -269,87 +269,87 @@ std::vector printFloatGenParameters = { // Default(right)-justified - { { "%f" }, "10.3456" }, + { "%f", "10.3456" }, // One position after the decimal,default(right)-justified - { { "%.1f" }, "10.3456" }, + { "%.1f", "10.3456" }, // Two positions after the decimal,default(right)-justified - { { "%.2f" }, "10.3456" }, + { "%.2f", "10.3456" }, //(Minimum)Eight-wide,three positions after the // decimal,default(right)-justified - { { "%8.3f" }, "10.3456" }, + { "%8.3f", "10.3456" }, //(Minimum)Eight-wide,two positions after the // decimal,zero-filled,default(right)-justified - { { "%08.2f" }, "10.3456" }, + { "%08.2f", "10.3456" }, //(Minimum)Eight-wide,two positions after the decimal,left-justified - { { "%-8.2f" }, "10.3456" }, + { "%-8.2f", "10.3456" }, //(Minimum)Eight-wide,two positions after the decimal,with // sign,default(right)-justified - { { "%+8.2f" }, "-10.3456" }, + { "%+8.2f", "-10.3456" }, // Zero positions after the // decimal([floor]rounding),default(right)-justified - { { "%.0f" }, "0.1" }, + { "%.0f", "0.1" }, // Zero positions after the decimal([ceil]rounding),default(right)-justified - { { "%.0f" }, "0.6" }, + { "%.0f", "0.6" }, // Zero-filled,default positions number after the // decimal,default(right)-justified - { { "%0f" }, "0.6" }, + { "%0f", "0.6" }, // Double argument representing floating-point,used by f // style,default(right)-justified - { { "%4g" }, "12345.6789" }, + { "%4g", "12345.6789" }, // Double argument representing floating-point,used by e // style,default(right)-justified - { { "%4.2g" }, "12345.6789" }, + { "%4.2g", "12345.6789" }, // Double argument representing floating-point,used by f // style,default(right)-justified - { { "%4G" }, "0.0000023" }, + { "%4G", "0.0000023" }, // Double argument representing floating-point,used by e // style,default(right)-justified - { { "%4G" }, "0.023" }, + { "%4G", "0.023" }, // Double argument representing floating-point,with // exponent,left-justified,default(right)-justified - { { "%-#20.15e" }, "789456123.0" }, + { "%-#20.15e", "789456123.0" }, // Double argument representing floating-point,with // exponent,left-justified,with sign,capital E,default(right)-justified - { { "%+#21.15E" }, "789456123.0" }, + { "%+#21.15E", "789456123.0" }, // Double argument representing floating-point,in [-]xh.hhhhpAd style - { { "%.6a" }, "0.1" }, + { "%.6a", "0.1" }, //(Minimum)Ten-wide,Double argument representing floating-point,in // xh.hhhhpAd style,default(right)-justified - { { "%10.2a" }, "9990.235" }, + { "%10.2a", "9990.235" }, }; //--------------------------------------------------------- @@ -391,14 +391,14 @@ std::vector printFloatLimitsGenParameters = { // Infinity (1.0/0.0) - { { "%f" }, "1.0f/0.0f" }, + { "%f", "1.0f/0.0f" }, // NaN - { { "%f" }, "sqrt(-1.0f)" }, + { "%f", "sqrt(-1.0f)" }, // NaN - { { "%f" }, "acospi(2.0f)" } + { "%f", "acospi(2.0f)" } }; //-------------------------------------------------------- @@ -453,21 +453,20 @@ std::vector printOctalGenParameters = { // Default(right)-justified - { { "%o" }, "10" }, + { "%o", "10" }, // Five-digit,default(right)-justified - { { "%.5o" }, "10" }, + { "%.5o", "10" }, // Default(right)-justified,increase precision - { { "%#o" }, "100000000" }, + { "%#o", "100000000" }, //(Minimum)Four-wide,Five-digit,0-flag ignored(because of // precision),default(right)-justified - { { "%04.5o" }, "10" } - + { "%04.5o", "10" } }; //------------------------------------------------------- @@ -510,11 +509,12 @@ std::vector printUnsignedGenParameters = { // Default(right)-justified - { { "%u" }, "10" }, + { "%u", "10" }, // Zero precision for zero,default(right)-justified - { { "%.0u" }, "0" }, + { "%.0u", "0" }, + }; @@ -558,23 +558,23 @@ std::vector printHexadecimalGenParameters = { // Add 0x,low x,default(right)-justified - { { "%#x" }, "0xABCDEF" }, + { "%#x", "0xABCDEF" }, // Add 0x,capital X,default(right)-justified - { { "%#X" }, "0xABCDEF" }, + { "%#X", "0xABCDEF" }, // Not add 0x,if zero,default(right)-justified - { { "%#X" }, "0" }, + { "%#X", "0" }, //(Minimum)Eight-wide,default(right)-justified - { { "%8x" }, "399" }, + { "%8x", "399" }, //(Minimum)Four-wide,zero-filled,default(right)-justified - { { "%04x" }, "399" } + { "%04x", "399" } }; @@ -618,15 +618,15 @@ std::vector printCharGenParameters = { // Four-wide,zero-filled,default(right)-justified - { { "%4c" }, "\'1\'" }, + { "%4c", "\'1\'" }, // Four-wide,left-justified - { { "%-4c" }, "\'1\'" }, + { "%-4c", "\'1\'" }, //(unsigned) int argument,default(right)-justified - { { "%c" }, "66" } + { "%c", "66" } }; @@ -690,36 +690,36 @@ std::vector printStringGenParameters = { //(Minimum)Four-wide,zero-filled,default(right)-justified - { { "%4s" }, "\"foo\"" }, + { "%4s", "\"foo\"" }, // One-digit(precision ignored),left-justified - { { "%.1s" }, "\"foo\"" }, + { "%.1s", "\"foo\"" }, //%% specification - { { "%s" }, "\"%%\"" }, + { "%s", "\"%%\"" }, // special symbols // nested - { { "%s" }, "\"\\\"%%\\\"\"" }, + { "%s", "\"\\\"%%\\\"\"" }, - { { "%s" }, "\"\\\'%%\\\'\"" }, + { "%s", "\"\\\'%%\\\'\"" }, // tabs - { { "%s" }, "\"foo\\tfoo\"" }, + { "%s", "\"foo\\tfoo\"" }, // newlines - { { "%s" }, "\"foo\\nfoo\"" }, + { "%s", "\"foo\\nfoo\"" }, // terminator - { { "%s" }, "\"foo\\0foo\"" }, + { "%s", "\"foo\\0foo\"" }, // all ascii characters - { { "%s" }, + { "%s", "\" " "!\\\"#$%&\'()*+,-./" "0123456789:;<=>?@ABCDEFGHIJKLMNOPQRSTUVWXYZ[\\\\]^_`" @@ -786,27 +786,27 @@ std::vector printFormatStringGenParameters = { //%% specification - { { "%%" } }, + { "%%" }, // special symbols // nested - { { "\\\"%%\\\"" } }, + { "\\\"%%\\\"" }, - { { "\'%%\'" } }, + { "\'%%\'" }, // tabs - { { "foo\\t\\t\\tfoo" } }, + { "foo\\t\\t\\tfoo" }, // newlines - { { "foo\\nfoo" } }, + { "foo\\nfoo" }, // all ascii characters - { { " !\\\"#$%&\'()*+,-./" + { " !\\\"#$%%&\'()*+,-./" "0123456789:;<=>?@ABCDEFGHIJKLMNOPQRSTUVWXYZ[\\\\]^_`" - "abcdefghijklmnopqrstuvwxyz{|}~" } } + "abcdefghijklmnopqrstuvwxyz{|}~" } }; //--------------------------------------------------------- @@ -873,27 +873,27 @@ std::vector printVectorGenParameters = { //(Minimum)Two-wide,two positions after decimal - { {}, "(1.0f,2.0f,3.0f,4.0f)", "%2.2", "hlf", "float", "4" }, + { NULL, "(1.0f,2.0f,3.0f,4.0f)", "%2.2", "hlf", "float", "4" }, // Alternative form,uchar argument - { {}, "(0xFA,0xFB)", "%#", "hhx", "uchar", "2" }, + { NULL, "(0xFA,0xFB)", "%#", "hhx", "uchar", "2" }, // Alternative form,ushort argument - { {}, "(0x1234,0x8765)", "%#", "hx", "ushort", "2" }, + { NULL, "(0x1234,0x8765)", "%#", "hx", "ushort", "2" }, // Alternative form,uint argument - { {}, "(0x12345678,0x87654321)", "%#", "hlx", "uint", "2" }, + { NULL, "(0x12345678,0x87654321)", "%#", "hlx", "uint", "2" }, // Alternative form,long argument - { {}, "(12345678,98765432)", "%", "ld", "long", "2" }, + { NULL, "(12345678,98765432)", "%", "ld", "long", "2" }, //(Minimum)Two-wide,two positions after decimal - { {}, "(1.0h,2.0h,3.0h,4.0h)", "%2.2", "hf", "half", "4" } + { NULL, "(1.0h,2.0h,3.0h,4.0h)", "%2.2", "hf", "half", "4" } }; //------------------------------------------------------------ @@ -957,67 +957,28 @@ std::vector printAddrSpaceGenParameters = { // Global memory region - { { "\"%d\\n\"" }, - NULL, - NULL, - NULL, - NULL, - NULL, - "__global int* x", - "", - "*x", + { "\"%d\\n\"", NULL, NULL, NULL, NULL, NULL, "__global int* x", "", "*x", "" }, // Global,constant, memory region - { { "\"%d\\n\"" }, - NULL, - NULL, - NULL, - NULL, - NULL, - "constant int* x", - "", - "*x", + { "\"%d\\n\"", NULL, NULL, NULL, NULL, NULL, "constant int* x", "", "*x", "" }, // Local memory region - { { "\"%+d\\n\"" }, - NULL, - NULL, - NULL, - NULL, - NULL, - "", - "local int x;\n x= (int)3;\n", - "x", - "" }, + { "\"%+d\\n\"", NULL, NULL, NULL, NULL, NULL, "", + "local int x;\n x= (int)3;\n", "x", "" }, // Private memory region - { { "\"%i\\n\"" }, - NULL, - NULL, - NULL, - NULL, - NULL, - "", - "private int x;\n x = (int)-1;\n", - "x", - "" }, + { "\"%i\\n\"", NULL, NULL, NULL, NULL, NULL, "", + "private int x;\n x = (int)-1;\n", "x", "" }, // Address of void * from global memory region - { { "\"%p\\n\"" }, - NULL, - NULL, - NULL, - NULL, - NULL, - "__global void* x,__global intptr_t* xAddr", - "", - "x", + { "\"%p\\n\"", NULL, NULL, NULL, NULL, NULL, + "__global void* x,__global intptr_t* xAddr", "", "x", "*xAddr = (intptr_t)x;\n" } }; @@ -1142,7 +1103,7 @@ size_t verifyOutputBuffer(char *analysisBuffer,testCase* pTestCase,size_t testId static void intRefBuilder(printDataGenParameters& params, char* refResult, const size_t refSize) { - snprintf(refResult, refSize, params.genericFormats.front().c_str(), + snprintf(refResult, refSize, params.genericFormat, atoi(params.dataRepresentation)); } @@ -1151,32 +1112,31 @@ static void halfRefBuilder(printDataGenParameters& params, char* refResult, { cl_half val = cl_half_from_float(strtof(params.dataRepresentation, NULL), half_rounding_mode); - snprintf(refResult, refSize, params.genericFormats.front().c_str(), - cl_half_to_float(val)); + snprintf(refResult, refSize, params.genericFormat, cl_half_to_float(val)); } static void floatRefBuilder(printDataGenParameters& params, char* refResult, const size_t refSize) { - snprintf(refResult, refSize, params.genericFormats.front().c_str(), + snprintf(refResult, refSize, params.genericFormat, strtof(params.dataRepresentation, NULL)); } static void octalRefBuilder(printDataGenParameters& params, char* refResult, const size_t refSize) { const unsigned long int data = strtoul(params.dataRepresentation, NULL, 10); - snprintf(refResult, refSize, params.genericFormats.front().c_str(), data); + snprintf(refResult, refSize, params.genericFormat, data); } static void unsignedRefBuilder(printDataGenParameters& params, char* refResult, const size_t refSize) { const unsigned long int data = strtoul(params.dataRepresentation, NULL, 10); - snprintf(refResult, refSize, params.genericFormats.front().c_str(), data); + snprintf(refResult, refSize, params.genericFormat, data); } static void hexRefBuilder(printDataGenParameters& params, char* refResult, const size_t refSize) { const unsigned long int data = strtoul(params.dataRepresentation, NULL, 0); - snprintf(refResult, refSize, params.genericFormats.front().c_str(), data); + snprintf(refResult, refSize, params.genericFormat, data); } /* From 02328fa38b5c5932d470381a257c58842e5c2556 Mon Sep 17 00:00:00 2001 From: Marcin Hajder Date: Tue, 28 May 2024 17:13:07 +0200 Subject: [PATCH 7/7] several corrections: -corrected bug in makePrintfProgram of using uninitialised name and extension -handled single printf argument use case -rearranged skip conditions in main doTest loop -clang format --- test_conformance/printf/test_printf.cpp | 535 ++++++++++++------------ 1 file changed, 277 insertions(+), 258 deletions(-) diff --git a/test_conformance/printf/test_printf.cpp b/test_conformance/printf/test_printf.cpp index fc6ef6c35..2ecf40018 100644 --- a/test_conformance/printf/test_printf.cpp +++ b/test_conformance/printf/test_printf.cpp @@ -231,69 +231,6 @@ static cl_program makePrintfProgram(cl_kernel* kernel_ptr, char addrSpacePAddArgument[256] = {0}; char extension[128] = { 0 }; - //Program Source code for int,float,octal,hexadecimal,char,string - const char* sourceGen[] = { - extension, - "__kernel void ", - testname, - "(void)\n", - "{\n" - " printf(\"", - allTestCase[testId] - ->_genParameters[testNum] - .genericFormats[formatNum] - .c_str(), - "\\n\",", - allTestCase[testId]->_genParameters[testNum].dataRepresentation, - ");", - "}\n" - }; - //Program Source code for vector - const char* sourceVec[] = { - extension, - "__kernel void ", - testname, - "(void)\n", - "{\n", - allTestCase[testId]->_genParameters[testNum].dataType, - allTestCase[testId]->_genParameters[testNum].vectorSize, - " tmp = (", - allTestCase[testId]->_genParameters[testNum].dataType, - allTestCase[testId]->_genParameters[testNum].vectorSize, - ")", - allTestCase[testId]->_genParameters[testNum].dataRepresentation, - ";", - " printf(\"", - allTestCase[testId]->_genParameters[testNum].vectorFormatFlag, - "v", - allTestCase[testId]->_genParameters[testNum].vectorSize, - allTestCase[testId]->_genParameters[testNum].vectorFormatSpecifier, - "\\n\",", - "tmp);", - "}\n" - }; - //Program Source code for address space - const char* sourceAddrSpace[] = { - "__kernel void ", - testname, - "(", - addrSpaceArgument, - ")\n{\n", - allTestCase[testId] - ->_genParameters[testNum] - .addrSpaceVariableTypeQualifier, - "printf(", - allTestCase[testId] - ->_genParameters[testNum] - .genericFormats[formatNum] - .c_str(), - ",", - allTestCase[testId]->_genParameters[testNum].addrSpaceParameter, - "); ", - addrSpacePAddArgument, - "\n}\n" - }; - //Update testname std::snprintf(testname, sizeof(testname), "%s%d", "test", testId); @@ -328,12 +265,59 @@ static cl_program makePrintfProgram(cl_kernel* kernel_ptr, strcpy(extension, "#pragma OPENCL EXTENSION cl_khr_fp16 : enable\n"); + // Program Source code for vector + const char* sourceVec[] = { + extension, + "__kernel void ", + testname, + "(void)\n", + "{\n", + allTestCase[testId]->_genParameters[testNum].dataType, + allTestCase[testId]->_genParameters[testNum].vectorSize, + " tmp = (", + allTestCase[testId]->_genParameters[testNum].dataType, + allTestCase[testId]->_genParameters[testNum].vectorSize, + ")", + allTestCase[testId]->_genParameters[testNum].dataRepresentation, + ";", + " printf(\"", + allTestCase[testId]->_genParameters[testNum].vectorFormatFlag, + "v", + allTestCase[testId]->_genParameters[testNum].vectorSize, + allTestCase[testId]->_genParameters[testNum].vectorFormatSpecifier, + "\\n\",", + "tmp);", + "}\n" + }; + err = create_single_kernel_helper( context, &program, kernel_ptr, sizeof(sourceVec) / sizeof(sourceVec[0]), sourceVec, testname); } else if(allTestCase[testId]->_type == TYPE_ADDRESS_SPACE) { + // Program Source code for address space + const char* sourceAddrSpace[] = { + "__kernel void ", + testname, + "(", + addrSpaceArgument, + ")\n{\n", + allTestCase[testId] + ->_genParameters[testNum] + .addrSpaceVariableTypeQualifier, + "printf(", + allTestCase[testId] + ->_genParameters[testNum] + .genericFormats[formatNum] + .c_str(), + ",", + allTestCase[testId]->_genParameters[testNum].addrSpaceParameter, + "); ", + addrSpacePAddArgument, + "\n}\n" + }; + err = create_single_kernel_helper(context, &program, kernel_ptr, sizeof(sourceAddrSpace) / sizeof(sourceAddrSpace[0]), @@ -341,9 +325,33 @@ static cl_program makePrintfProgram(cl_kernel* kernel_ptr, } else { - err = create_single_kernel_helper( - context, &program, kernel_ptr, - sizeof(sourceGen) / sizeof(sourceGen[0]), sourceGen, testname); + // Program Source code for int,float,octal,hexadecimal,char,string + std::ostringstream sourceGen; + sourceGen << extension << "__kernel void " << testname + << "(void)\n" + "{\n" + " printf(\"" + << allTestCase[testId] + ->_genParameters[testNum] + .genericFormats[formatNum] + .c_str() + << "\\n\""; + + if (allTestCase[testId]->_genParameters[testNum].dataRepresentation) + { + sourceGen << "," + << allTestCase[testId] + ->_genParameters[testNum] + .dataRepresentation; + } + + sourceGen << ");\n}\n"; + + std::string kernel_source = sourceGen.str(); + const char* ptr = kernel_source.c_str(); + + err = create_single_kernel_helper(context, &program, kernel_ptr, 1, + &ptr, testname); } if (!program || err) { @@ -465,13 +473,6 @@ static int doTest(cl_command_queue queue, cl_context context, for (unsigned testNum = 0; testNum < genParams.size(); testNum++) { - - - for (unsigned formatNum = 0; formatNum - < allTestCase[testId]->_genParameters[testNum].genericFormats.size(); - formatNum++) - { - if (allTestCase[testId]->_type == TYPE_VECTOR) { if ((strcmp(allTestCase[testId]->_genParameters[testNum].dataType, @@ -487,227 +488,245 @@ static int doTest(cl_command_queue queue, cl_context context, continue; } - 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); + // Long support for varible type + if (!strcmp(allTestCase[testId]->_genParameters[testNum].dataType, + "long") + && !isLongSupported(device)) + { + log_info("Long is not supported, test not run.\n"); + s_test_skip++; + s_test_cnt++; + continue; + } } - else if (allTestCase[testId]->_type == TYPE_ADDRESS_SPACE) + + for (unsigned formatNum = 0; formatNum < allTestCase[testId] + ->_genParameters[testNum] + .genericFormats.size(); + formatNum++) { - if (isKernelArgument(allTestCase[testId], testNum)) + if (allTestCase[testId]->_type == TYPE_VECTOR) { - 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); + 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 kernel //variable %s \n printf(%s,%s)\n", - testNum, - allTestCase[testId] - ->_genParameters[testNum] - .addrSpaceVariableTypeQualifier, + log_info("%d)testing printf(\"%s\"", testNum, allTestCase[testId] ->_genParameters[testNum] .genericFormats[formatNum] - .c_str(), - allTestCase[testId] - ->_genParameters[testNum] - .addrSpaceParameter); + .c_str()); + if (allTestCase[testId] + ->_genParameters[testNum] + .dataRepresentation) + log_info(",%s", + allTestCase[testId] + ->_genParameters[testNum] + .dataRepresentation); + log_info(")\n"); } - } - else - { - log_info("%d)testing printf(\"%s\",%s)\n", testNum, - allTestCase[testId] - ->_genParameters[testNum] - .genericFormats[formatNum] - .c_str(), - allTestCase[testId] - ->_genParameters[testNum] - .dataRepresentation); - } - - // Long support for varible type - if (allTestCase[testId]->_type == TYPE_VECTOR - && !strcmp(allTestCase[testId]->_genParameters[testNum].dataType, - "long") - && !isLongSupported(device)) - { - log_info("Long is not supported, test not run.\n"); - s_test_skip++; - s_test_cnt++; - continue; - } - fflush(stdout); + 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; - } + // 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; + } - clProgramWrapper program; - clKernelWrapper kernel; - clMemWrapper d_out; - clMemWrapper d_a; - char _analysisBuffer[ANALYSIS_BUFFER_SIZE]; - cl_uint out32 = 0; - cl_ulong out64 = 0; - int fd = -1; - - // Define an index space (global work size) of threads for execution. - size_t globalWorkSize[1]; - - program = - makePrintfProgram(&kernel, context, testId, testNum, formatNum, - isLongSupport, is64bAddressSpace(device)); - if (!program || !kernel) - { - subtest_fail(nullptr); - continue; - } + clProgramWrapper program; + clKernelWrapper kernel; + clMemWrapper d_out; + clMemWrapper d_a; + char _analysisBuffer[ANALYSIS_BUFFER_SIZE]; + cl_uint out32 = 0; + cl_ulong out64 = 0; + int fd = -1; + + // Define an index space (global work size) of threads for + // execution. + size_t globalWorkSize[1]; + + program = + makePrintfProgram(&kernel, context, testId, testNum, formatNum, + isLongSupport, is64bAddressSpace(device)); + if (!program || !kernel) + { + subtest_fail(nullptr); + continue; + } - // For address space test if there is kernel argument - set it - if (allTestCase[testId]->_type == TYPE_ADDRESS_SPACE) - { - if (isKernelArgument(allTestCase[testId], testNum)) + // For address space test if there is kernel argument - set it + if (allTestCase[testId]->_type == TYPE_ADDRESS_SPACE) { - int a = 2; - d_a = clCreateBuffer(context, - CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, - sizeof(int), &a, &err); - if (err != CL_SUCCESS || d_a == NULL) + if (isKernelArgument(allTestCase[testId], testNum)) { - subtest_fail("clCreateBuffer failed\n"); - continue; + int a = 2; + d_a = clCreateBuffer( + context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, + sizeof(int), &a, &err); + if (err != CL_SUCCESS || d_a == NULL) + { + subtest_fail("clCreateBuffer failed\n"); + continue; + } + err = clSetKernelArg(kernel, 0, sizeof(cl_mem), &d_a); + if (err != CL_SUCCESS) + { + subtest_fail("clSetKernelArg failed\n"); + continue; + } } - err = clSetKernelArg(kernel, 0, sizeof(cl_mem), &d_a); - if (err != CL_SUCCESS) + // For address space test if %p is tested + if (isKernelPFormat(allTestCase[testId], testNum)) { - subtest_fail("clSetKernelArg failed\n"); - continue; + d_out = clCreateBuffer(context, CL_MEM_READ_WRITE, + sizeof(cl_ulong), NULL, &err); + if (err != CL_SUCCESS || d_out == NULL) + { + subtest_fail("clCreateBuffer failed\n"); + continue; + } + err = clSetKernelArg(kernel, 1, sizeof(cl_mem), &d_out); + if (err != CL_SUCCESS) + { + subtest_fail("clSetKernelArg failed\n"); + continue; + } } } - // For address space test if %p is tested - if (isKernelPFormat(allTestCase[testId], testNum)) + + fd = acquireOutputStream(&err); + if (err != 0) { - d_out = clCreateBuffer(context, CL_MEM_READ_WRITE, - sizeof(cl_ulong), NULL, &err); - if (err != CL_SUCCESS || d_out == NULL) - { - subtest_fail("clCreateBuffer failed\n"); - continue; - } - err = clSetKernelArg(kernel, 1, sizeof(cl_mem), &d_out); - if (err != CL_SUCCESS) - { - subtest_fail("clSetKernelArg failed\n"); - continue; - } + subtest_fail("Error while redirection stdout to file"); + continue; + } + globalWorkSize[0] = 1; + cl_event ndrEvt; + err = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, globalWorkSize, + NULL, 0, NULL, &ndrEvt); + if (err != CL_SUCCESS) + { + releaseOutputStream(fd); + subtest_fail("\n clEnqueueNDRangeKernel failed errcode:%d\n", + err); + continue; } - } - fd = acquireOutputStream(&err); - if (err != 0) - { - subtest_fail("Error while redirection stdout to file"); - continue; - } - globalWorkSize[0] = 1; - cl_event ndrEvt; - err = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, globalWorkSize, - NULL, 0, NULL, &ndrEvt); - if (err != CL_SUCCESS) - { - releaseOutputStream(fd); - subtest_fail("\n clEnqueueNDRangeKernel failed errcode:%d\n", err); - continue; - } + fflush(stdout); + err = clFlush(queue); + if (err != CL_SUCCESS) + { + releaseOutputStream(fd); + subtest_fail("clFlush failed : %d\n", err); + continue; + } + // Wait until kernel finishes its execution and (thus) the output + // printed from the kernel is immediately printed + err = waitForEvent(&ndrEvt); - fflush(stdout); - err = clFlush(queue); - if (err != CL_SUCCESS) - { releaseOutputStream(fd); - subtest_fail("clFlush failed : %d\n", err); - continue; - } - // Wait until kernel finishes its execution and (thus) the output - // printed from the kernel is immediately printed - err = waitForEvent(&ndrEvt); - releaseOutputStream(fd); - - if (err != CL_SUCCESS) - { - subtest_fail("waitforEvent failed : %d\n", err); - continue; - } - fflush(stdout); - - if (allTestCase[testId]->_type == TYPE_ADDRESS_SPACE - && isKernelPFormat(allTestCase[testId], testNum)) - { - // Read the OpenCL output buffer (d_out) to the host output array - // (out) - if (!is64bAddressSpace(device)) // 32-bit address space + if (err != CL_SUCCESS) { - clEnqueueReadBuffer(queue, d_out, CL_TRUE, 0, sizeof(cl_int), - &out32, 0, NULL, NULL); + subtest_fail("waitforEvent failed : %d\n", err); + continue; } - else // 64-bit address space + fflush(stdout); + + if (allTestCase[testId]->_type == TYPE_ADDRESS_SPACE + && isKernelPFormat(allTestCase[testId], testNum)) { - clEnqueueReadBuffer(queue, d_out, CL_TRUE, 0, sizeof(cl_ulong), - &out64, 0, NULL, NULL); + // Read the OpenCL output buffer (d_out) to the host output + // array (out) + if (!is64bAddressSpace(device)) // 32-bit address space + { + clEnqueueReadBuffer(queue, d_out, CL_TRUE, 0, + sizeof(cl_int), &out32, 0, NULL, NULL); + } + else // 64-bit address space + { + clEnqueueReadBuffer(queue, d_out, CL_TRUE, 0, + sizeof(cl_ulong), &out64, 0, NULL, + NULL); + } } - } - // - // Get the output printed from the kernel to _analysisBuffer - // and verify its correctness - getAnalysisBuffer(_analysisBuffer); - if (!is64bAddressSpace(device)) // 32-bit address space - { - if (0 - != verifyOutputBuffer(_analysisBuffer, allTestCase[testId], - testNum, (cl_ulong)out32)) + // + // Get the output printed from the kernel to _analysisBuffer + // and verify its correctness + getAnalysisBuffer(_analysisBuffer); + if (!is64bAddressSpace(device)) // 32-bit address space { - subtest_fail("verifyOutputBuffer failed\n"); - continue; + if (0 + != verifyOutputBuffer(_analysisBuffer, allTestCase[testId], + testNum, (cl_ulong)out32)) + { + subtest_fail("verifyOutputBuffer failed\n"); + continue; + } } - } - else //64-bit address space - { - if (0 - != verifyOutputBuffer(_analysisBuffer, allTestCase[testId], - testNum, out64)) + else // 64-bit address space { - subtest_fail("verifyOutputBuffer failed\n"); - continue; + if (0 + != verifyOutputBuffer(_analysisBuffer, allTestCase[testId], + testNum, out64)) + { + subtest_fail("verifyOutputBuffer failed\n"); + continue; + } } } - - } ++s_test_cnt; }