From e938dd8cf8b91cb82d95e2937067ede3ccd177d4 Mon Sep 17 00:00:00 2001 From: Marcin Hajder Date: Wed, 19 Jun 2024 08:35:05 +0200 Subject: [PATCH 1/5] Added mixed format printf subtest --- test_conformance/printf/test_printf.cpp | 445 ++++++++++++++++-------- test_conformance/printf/test_printf.h | 1 + test_conformance/printf/util_printf.cpp | 66 +++- 3 files changed, 364 insertions(+), 148 deletions(-) diff --git a/test_conformance/printf/test_printf.cpp b/test_conformance/printf/test_printf.cpp index 2ecf40018..0fe517e14 100644 --- a/test_conformance/printf/test_printf.cpp +++ b/test_conformance/printf/test_printf.cpp @@ -15,7 +15,9 @@ // #include "harness/os_helpers.h" #include "harness/typeWrappers.h" +#include "harness/stringHelpers.h" +#include #include #include #include @@ -43,6 +45,7 @@ #include "harness/errorHelpers.h" #include "harness/kernelHelpers.h" #include "harness/parseParameters.h" +#include "harness/rounding_mode.h" #include @@ -51,50 +54,49 @@ typedef unsigned int uint32_t; test_status InitCL( cl_device_id device ); +namespace { + //----------------------------------------- -// Static helper functions declaration +// helper functions declaration //----------------------------------------- -static void printUsage( void ); - //Stream helper functions //Associate stdout stream with the file(gFileName):i.e redirect stdout stream to the specific files (gFileName) -static int acquireOutputStream(int* error); +int acquireOutputStream(int* error); //Close the file(gFileName) associated with the stdout stream and disassociates it. -static void releaseOutputStream(int fd); +void releaseOutputStream(int fd); //Get analysis buffer to verify the correctess of printed data -static void getAnalysisBuffer(char* analysisBuffer); +void getAnalysisBuffer(char* analysisBuffer); //Kernel builder helper functions //Check if the test case is for kernel that has argument -static int isKernelArgument(testCase* pTestCase,size_t testId); +int isKernelArgument(testCase* pTestCase, size_t testId); //Check if the test case treats %p format for void* -static int isKernelPFormat(testCase* pTestCase,size_t testId); +int isKernelPFormat(testCase* pTestCase, size_t testId); //----------------------------------------- // Static functions declarations //----------------------------------------- // Make a program that uses printf for the given type/format, -static cl_program -makePrintfProgram(cl_kernel* kernel_ptr, const cl_context context, - const unsigned int testId, const unsigned int testNum, - const unsigned int formatNum, bool isLongSupport = true, - bool is64bAddrSpace = false); +cl_program makePrintfProgram(cl_kernel* kernel_ptr, const cl_context context, + cl_device_id device, const unsigned int testId, + const unsigned int testNum, + const unsigned int formatNum); // Creates and execute the printf test for the given device, context, type/format -static int doTest(cl_command_queue queue, cl_context context, - const unsigned int testId, cl_device_id device); +int doTest(cl_command_queue queue, cl_context context, + const unsigned int testId, cl_device_id device); // Check if device supports long -static bool isLongSupported(cl_device_id device_id); +bool isLongSupported(cl_device_id device_id); // Check if device address space is 64 bits -static bool is64bAddressSpace(cl_device_id device_id); +bool is64bAddressSpace(cl_device_id device_id); //Wait until event status is CL_COMPLETE int waitForEvent(cl_event* event); @@ -111,21 +113,25 @@ int s_test_cnt = 0; int s_test_fail = 0; int s_test_skip = 0; +cl_context gContext; +cl_command_queue gQueue; +int gFd; + +char gFileName[256]; -static cl_context gContext; -static cl_command_queue gQueue; -static int gFd; +MTdataHolder gMTdata; -static char gFileName[256]; +// For the sake of proper logging of negative results +std::string gLatestKernelSource; //----------------------------------------- -// Static helper functions definition +// helper functions definition //----------------------------------------- //----------------------------------------- // acquireOutputStream //----------------------------------------- -static int acquireOutputStream(int* error) +int acquireOutputStream(int* error) { int fd = streamDup(fileno(stdout)); *error = 0; @@ -140,7 +146,7 @@ static int acquireOutputStream(int* error) //----------------------------------------- // releaseOutputStream //----------------------------------------- -static void releaseOutputStream(int fd) +void releaseOutputStream(int fd) { fflush(stdout); streamDup2(fd,fileno(stdout)); @@ -150,7 +156,8 @@ static void releaseOutputStream(int fd) //----------------------------------------- // printfCallBack //----------------------------------------- -static void CL_CALLBACK printfCallBack(const char *printf_data, size_t len, size_t final, void *user_data) +void CL_CALLBACK printfCallBack(const char* printf_data, size_t len, + size_t final, void* user_data) { fwrite(printf_data, 1, len, stdout); } @@ -158,7 +165,7 @@ static void CL_CALLBACK printfCallBack(const char *printf_data, size_t len, size //----------------------------------------- // getAnalysisBuffer //----------------------------------------- -static void getAnalysisBuffer(char* analysisBuffer) +void getAnalysisBuffer(char* analysisBuffer) { FILE *fp; memset(analysisBuffer,0,ANALYSIS_BUFFER_SIZE); @@ -177,14 +184,14 @@ static void getAnalysisBuffer(char* analysisBuffer) //----------------------------------------- // isKernelArgument //----------------------------------------- -static int isKernelArgument(testCase* pTestCase,size_t testId) +int isKernelArgument(testCase* pTestCase, size_t testId) { return strcmp(pTestCase->_genParameters[testId].addrSpaceArgumentTypeQualifier,""); } //----------------------------------------- // isKernelPFormat //----------------------------------------- -static int isKernelPFormat(testCase* pTestCase,size_t testId) +int isKernelPFormat(testCase* pTestCase, size_t testId) { return strcmp(pTestCase->_genParameters[testId].addrSpacePAdd,""); } @@ -214,15 +221,148 @@ int waitForEvent(cl_event* event) // Static helper functions definition //----------------------------------------- +float get_random_float(float low, float high, MTdata d) +{ + float t = (float)((double)genrand_int32(d) / (double)0xFFFFFFFF); + return (1.0f - t) * low + t * high; +} + +//----------------------------------------- +// makeMixedFormatPrintfProgram +//----------------------------------------- +cl_program makeMixedFormatPrintfProgram(cl_kernel* kernel_ptr, + const cl_context context, + const cl_device_id device, + const unsigned int testId, + const std::string& testname) +{ + auto gen_char = [&]() { + static const char dict[] = { + " !#$&()*+,-./" + "123456789:;<=>?@ABCDEFGHIJKLMNOPQRSTUVWXYZ[]^_`" + "abcdefghijklmnopqrstuvwxyz{|}~" + }; + return dict[genrand_int32(gMTdata) % ((int)sizeof(dict) - 1)]; + }; + + std::array, 2> formats = { + { { "%f", "%e", "%g", "%a", "%F", "%E", "%G", "%A" }, + { "%d", "%i", "%u", "%x" } } + }; + std::vector data_befor(2 + genrand_int32(gMTdata) % 8); + std::vector data_after(2 + genrand_int32(gMTdata) % 8); + + std::generate(data_befor.begin(), data_befor.end(), gen_char); + std::generate(data_after.begin(), data_after.end(), gen_char); + + cl_uint num_args = 2 + genrand_int32(gMTdata) % 4; + + // Map device rounding to CTS rounding type + // get_default_rounding_mode supports RNE and RTZ + auto get_rounding = [](const cl_device_fp_config& fpConfig) { + if (fpConfig == CL_FP_ROUND_TO_NEAREST) + { + return kRoundToNearestEven; + } + else if (fpConfig == CL_FP_ROUND_TO_ZERO) + { + return kRoundTowardZero; + } + else + { + assert(false && "Unreachable"); + } + return kDefaultRoundingMode; + }; + + const RoundingMode hostRound = get_round(); + RoundingMode deviceRound = get_rounding(get_default_rounding_mode(device)); + + std::ostringstream format_str; + std::ostringstream ref_str; + std::ostringstream source_gen; + std::ostringstream args_str; + source_gen << "__kernel void " << testname + << "(void)\n" + "{\n" + " printf(\""; + for (auto it : data_befor) + { + format_str << it; + ref_str << it; + } + format_str << ", "; + ref_str << ", "; + + + for (cl_uint i = 0; i < num_args; i++) + { + std::uint8_t type_ind = genrand_int32(gMTdata) % 2; + + // Set CPU rounding mode to match that of the device + set_round(deviceRound, type_ind != 0 ? kint : kfloat); + + std::string format = formats[type_ind][genrand_int32(gMTdata) + % formats[type_ind].size()]; + format_str << format << ", "; + + if (type_ind) + { + int arg = genrand_int32(gMTdata); + args_str << str_sprintf("%d", arg) << ", "; + ref_str << str_sprintf(format, arg) << ", "; + } + else + { + const float max_range = 100000.f; + float arg = get_random_float(-max_range, max_range, gMTdata); + args_str << str_sprintf("%f", arg) << "f, "; + ref_str << str_sprintf(format, arg) << ", "; + } + } + // Restore the original CPU rounding mode + set_round(hostRound, kfloat); + + args_str.seekp(-2, std::ios_base::end); + args_str << ");\n}\n"; + + for (auto it : data_after) + { + format_str << it; + ref_str << it; + } + source_gen << format_str.str() << "\\n\"" + << ", " << args_str.str(); + + std::string kernel_source = source_gen.str(); + const char* ptr = kernel_source.c_str(); + + cl_program program; + cl_int err = create_single_kernel_helper(context, &program, kernel_ptr, 1, + &ptr, testname.c_str()); + + gLatestKernelSource = kernel_source.c_str(); + + // Save the reference result + allTestCase[testId]->_correctBuffer.push_back(ref_str.str()); + + if (!program || err) + { + log_error("create_single_kernel_helper failed\n"); + return NULL; + } + + return program; +} + //----------------------------------------- // makePrintfProgram //----------------------------------------- -static cl_program makePrintfProgram(cl_kernel* kernel_ptr, - const cl_context context, - const unsigned int testId, - const unsigned int testNum, - const unsigned int formatNum, - bool isLongSupport, bool is64bAddrSpace) +cl_program makePrintfProgram(cl_kernel* kernel_ptr, const cl_context context, + const cl_device_id device, + const unsigned int testId, + const unsigned int testNum, + const unsigned int formatNum) { int err; cl_program program; @@ -293,6 +433,9 @@ static cl_program makePrintfProgram(cl_kernel* kernel_ptr, err = create_single_kernel_helper( context, &program, kernel_ptr, sizeof(sourceVec) / sizeof(sourceVec[0]), sourceVec, testname); + + gLatestKernelSource = + concat_kernel(sourceVec, sizeof(sourceVec) / sizeof(sourceVec[0])); } else if(allTestCase[testId]->_type == TYPE_ADDRESS_SPACE) { @@ -322,6 +465,15 @@ static cl_program makePrintfProgram(cl_kernel* kernel_ptr, sizeof(sourceAddrSpace) / sizeof(sourceAddrSpace[0]), sourceAddrSpace, testname); + + gLatestKernelSource = + concat_kernel(sourceAddrSpace, + sizeof(sourceAddrSpace) / sizeof(sourceAddrSpace[0])); + } + else if (allTestCase[testId]->_type == TYPE_MIXED_FORMAT) + { + return makeMixedFormatPrintfProgram(kernel_ptr, context, device, testId, + testname); } else { @@ -352,6 +504,8 @@ static cl_program makePrintfProgram(cl_kernel* kernel_ptr, err = create_single_kernel_helper(context, &program, kernel_ptr, 1, &ptr, testname); + + gLatestKernelSource = kernel_source.c_str(); } if (!program || err) { @@ -365,7 +519,7 @@ static cl_program makePrintfProgram(cl_kernel* kernel_ptr, //----------------------------------------- // isLongSupported //----------------------------------------- -static bool isLongSupported(cl_device_id device_id) +bool isLongSupported(cl_device_id device_id) { size_t tempSize = 0; cl_int status; @@ -409,7 +563,7 @@ static bool isLongSupported(cl_device_id device_id) //----------------------------------------- // is64bAddressSpace //----------------------------------------- -static bool is64bAddressSpace(cl_device_id device_id) +bool is64bAddressSpace(cl_device_id device_id) { cl_int status; cl_uint addrSpaceB; @@ -448,11 +602,82 @@ void subtest_fail(const char* msg, ...) ++s_test_cnt; } +//----------------------------------------- +// logTestType - printout test details +//----------------------------------------- + +void logTestType(const unsigned testId, const unsigned testNum, + unsigned formatNum) +{ + if (allTestCase[testId]->_type == TYPE_VECTOR) + { + log_info( + "%d)testing printf(\"%sv%s%s\",%s)\n", testNum, + allTestCase[testId]->_genParameters[testNum].vectorFormatFlag, + allTestCase[testId]->_genParameters[testNum].vectorSize, + allTestCase[testId]->_genParameters[testNum].vectorFormatSpecifier, + allTestCase[testId]->_genParameters[testNum].dataRepresentation); + } + else if (allTestCase[testId]->_type == TYPE_ADDRESS_SPACE) + { + if (isKernelArgument(allTestCase[testId], testNum)) + { + log_info("%d)testing kernel //argument %s \n printf(%s,%s)\n", + testNum, + allTestCase[testId] + ->_genParameters[testNum] + .addrSpaceArgumentTypeQualifier, + allTestCase[testId] + ->_genParameters[testNum] + .genericFormats[formatNum] + .c_str(), + allTestCase[testId] + ->_genParameters[testNum] + .addrSpaceParameter); + } + else + { + log_info("%d)testing kernel //variable %s \n printf(%s,%s)\n", + testNum, + allTestCase[testId] + ->_genParameters[testNum] + .addrSpaceVariableTypeQualifier, + allTestCase[testId] + ->_genParameters[testNum] + .genericFormats[formatNum] + .c_str(), + allTestCase[testId] + ->_genParameters[testNum] + .addrSpaceParameter); + } + } + else if (allTestCase[testId]->_type == TYPE_MIXED_FORMAT) + { + log_info("%d) testing printf with mixed format string\n", testNum); + } + else + { + log_info("%d)testing printf(\"%s\"", testNum, + allTestCase[testId] + ->_genParameters[testNum] + .genericFormats[formatNum] + .c_str()); + if (allTestCase[testId]->_genParameters[testNum].dataRepresentation) + log_info(",%s", + allTestCase[testId] + ->_genParameters[testNum] + .dataRepresentation); + log_info(")\n"); + } + + fflush(stdout); +} + //----------------------------------------- // doTest //----------------------------------------- -static int doTest(cl_command_queue queue, cl_context context, - const unsigned int testId, cl_device_id device) +int doTest(cl_command_queue queue, cl_context context, + const unsigned int testId, cl_device_id device) { int err = TEST_FAIL; @@ -500,88 +725,13 @@ static int doTest(cl_command_queue queue, cl_context context, } } - for (unsigned formatNum = 0; formatNum < allTestCase[testId] - ->_genParameters[testNum] - .genericFormats.size(); + auto genParamsVec = allTestCase[testId]->_genParameters; + auto genFormatVec = genParamsVec[testNum].genericFormats; + + for (unsigned formatNum = 0; formatNum < genFormatVec.size(); formatNum++) { - if (allTestCase[testId]->_type == TYPE_VECTOR) - { - log_info( - "%d)testing printf(\"%sv%s%s\",%s)\n", testNum, - allTestCase[testId] - ->_genParameters[testNum] - .vectorFormatFlag, - allTestCase[testId]->_genParameters[testNum].vectorSize, - allTestCase[testId] - ->_genParameters[testNum] - .vectorFormatSpecifier, - allTestCase[testId] - ->_genParameters[testNum] - .dataRepresentation); - } - else if (allTestCase[testId]->_type == TYPE_ADDRESS_SPACE) - { - if (isKernelArgument(allTestCase[testId], testNum)) - { - log_info( - "%d)testing kernel //argument %s \n printf(%s,%s)\n", - testNum, - allTestCase[testId] - ->_genParameters[testNum] - .addrSpaceArgumentTypeQualifier, - allTestCase[testId] - ->_genParameters[testNum] - .genericFormats[formatNum] - .c_str(), - allTestCase[testId] - ->_genParameters[testNum] - .addrSpaceParameter); - } - else - { - log_info( - "%d)testing kernel //variable %s \n printf(%s,%s)\n", - testNum, - allTestCase[testId] - ->_genParameters[testNum] - .addrSpaceVariableTypeQualifier, - allTestCase[testId] - ->_genParameters[testNum] - .genericFormats[formatNum] - .c_str(), - allTestCase[testId] - ->_genParameters[testNum] - .addrSpaceParameter); - } - } - else - { - log_info("%d)testing printf(\"%s\"", testNum, - allTestCase[testId] - ->_genParameters[testNum] - .genericFormats[formatNum] - .c_str()); - if (allTestCase[testId] - ->_genParameters[testNum] - .dataRepresentation) - log_info(",%s", - allTestCase[testId] - ->_genParameters[testNum] - .dataRepresentation); - log_info(")\n"); - } - - fflush(stdout); - - // Long support for address in FULL_PROFILE/EMBEDDED_PROFILE - bool isLongSupport = true; - if (allTestCase[testId]->_type == TYPE_ADDRESS_SPACE - && isKernelPFormat(allTestCase[testId], testNum) - && !isLongSupported(device)) - { - isLongSupport = false; - } + logTestType(testId, testNum, formatNum); clProgramWrapper program; clKernelWrapper kernel; @@ -596,9 +746,8 @@ static int doTest(cl_command_queue queue, cl_context context, // execution. size_t globalWorkSize[1]; - program = - makePrintfProgram(&kernel, context, testId, testNum, formatNum, - isLongSupport, is64bAddressSpace(device)); + program = makePrintfProgram(&kernel, context, device, testId, + testNum, formatNum); if (!program || !kernel) { subtest_fail(nullptr); @@ -712,7 +861,12 @@ static int doTest(cl_command_queue queue, cl_context context, != verifyOutputBuffer(_analysisBuffer, allTestCase[testId], testNum, (cl_ulong)out32)) { - subtest_fail("verifyOutputBuffer failed\n"); + subtest_fail( + "verifyOutputBuffer failed with kernel: " + "\n%s\n expected: %s\n got: %s\n", + gLatestKernelSource.c_str(), + allTestCase[testId]->_correctBuffer[testNum].c_str(), + _analysisBuffer); continue; } } @@ -722,7 +876,12 @@ static int doTest(cl_command_queue queue, cl_context context, != verifyOutputBuffer(_analysisBuffer, allTestCase[testId], testNum, out64)) { - subtest_fail("verifyOutputBuffer failed\n"); + subtest_fail( + "verifyOutputBuffer failed with kernel: " + "\n%s\n expected: %s\n got: %s\n", + gLatestKernelSource.c_str(), + allTestCase[testId]->_correctBuffer[testNum].c_str(), + _analysisBuffer); continue; } } @@ -736,6 +895,8 @@ static int doTest(cl_command_queue queue, cl_context context, return s_test_fail - fail_count; } +} + int test_int(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) { @@ -814,6 +975,12 @@ int test_address_space(cl_device_id deviceID, cl_context context, return doTest(gQueue, gContext, TYPE_ADDRESS_SPACE, deviceID); } +int test_mixed_format(cl_device_id deviceID, cl_context context, + cl_command_queue queue, int num_elements) +{ + return doTest(gQueue, gContext, TYPE_MIXED_FORMAT, deviceID); +} + int test_buffer_size(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) { @@ -844,11 +1011,25 @@ test_definition test_list[] = { ADD_TEST(float), ADD_TEST(float_limits), ADD_TEST(octal), ADD_TEST(unsigned), ADD_TEST(hexadecimal), ADD_TEST(char), ADD_TEST(string), ADD_TEST(format_string), ADD_TEST(vector), - ADD_TEST(address_space), ADD_TEST(buffer_size), + ADD_TEST(address_space), ADD_TEST(buffer_size), ADD_TEST(mixed_format), }; const int test_num = ARRAY_SIZE( test_list ); +//----------------------------------------- +// printUsage +//----------------------------------------- +static void printUsage(void) +{ + log_info("test_printf: \n"); + log_info("\tdefault is to run the full test on the default device\n"); + log_info("\n"); + for (int i = 0; i < test_num; i++) + { + log_info("\t%s\n", test_list[i].name); + } +} + //----------------------------------------- // main //----------------------------------------- @@ -913,6 +1094,8 @@ int main(int argc, const char* argv[]) return -1; } + gMTdata = MTdataHolder(gRandomSeed); + int err = runTestHarnessWithCheck( argCount, argList, test_num, test_list, true, 0, InitCL ); if(gQueue) @@ -934,20 +1117,6 @@ int main(int argc, const char* argv[]) return err; } -//----------------------------------------- -// printUsage -//----------------------------------------- -static void printUsage( void ) -{ - log_info("test_printf: \n"); - log_info("\tdefault is to run the full test on the default device\n"); - log_info("\n"); - for( int i = 0; i < test_num; i++ ) - { - log_info( "\t%s\n", test_list[i].name ); - } -} - test_status InitCL( cl_device_id device ) { uint32_t device_frequency = 0; diff --git a/test_conformance/printf/test_printf.h b/test_conformance/printf/test_printf.h index 0a33d5f84..f2d667172 100644 --- a/test_conformance/printf/test_printf.h +++ b/test_conformance/printf/test_printf.h @@ -58,6 +58,7 @@ enum PrintfTestType TYPE_FORMAT_STRING, TYPE_VECTOR, TYPE_ADDRESS_SPACE, + TYPE_MIXED_FORMAT, TYPE_COUNT }; diff --git a/test_conformance/printf/util_printf.cpp b/test_conformance/printf/util_printf.cpp index ca260573b..7ce76dd75 100644 --- a/test_conformance/printf/util_printf.cpp +++ b/test_conformance/printf/util_printf.cpp @@ -1078,7 +1078,24 @@ testCase testCaseAddrSpace = { }; +//========================================================= +// mixed format +//========================================================= +//---------------------------------------------------------- +// placeholders for mixed-args-data test +//---------------------------------------------------------- +// empty slots specifies number of tests +std::vector printMixedFormatGenParameters(64, + { { "" } }); + +std::vector correctBufferMixedFormat; + +//---------------------------------------------------------- +// Test case for mixed-args +//---------------------------------------------------------- +testCase testCaseMixedFormat = { TYPE_MIXED_FORMAT, correctBufferMixedFormat, + printMixedFormatGenParameters, NULL }; //------------------------------------------------------------------------------- @@ -1087,11 +1104,11 @@ testCase testCaseAddrSpace = { //------------------------------------------------------------------------------- std::vector allTestCase = { - &testCaseInt, &testCaseHalf, &testCaseHalfLimits, - &testCaseFloat, &testCaseFloatLimits, &testCaseOctal, - &testCaseUnsigned, &testCaseHexadecimal, &testCaseChar, - &testCaseString, &testCaseFormatString, &testCaseVector, - &testCaseAddrSpace + &testCaseInt, &testCaseHalf, &testCaseHalfLimits, + &testCaseFloat, &testCaseFloatLimits, &testCaseOctal, + &testCaseUnsigned, &testCaseHexadecimal, &testCaseChar, + &testCaseString, &testCaseFormatString, &testCaseVector, + &testCaseAddrSpace, &testCaseMixedFormat }; //----------------------------------------- @@ -1134,14 +1151,29 @@ size_t verifyOutputBuffer(char *analysisBuffer,testCase* pTestCase,size_t testId } - char* exp; - //Exponenent representation - if((exp = strstr(analysisBuffer,"E+")) != NULL || (exp = strstr(analysisBuffer,"e+")) != NULL || (exp = strstr(analysisBuffer,"E-")) != NULL || (exp = strstr(analysisBuffer,"e-")) != NULL) + char* exp = nullptr; + std::string copy_str; + std::vector staging(strlen(analysisBuffer) + 1); + std::vector staging_correct(pTestCase->_correctBuffer[testId].size() + + 1); + std::snprintf(staging.data(), staging.size(), "%s", analysisBuffer); + std::snprintf(staging_correct.data(), staging_correct.size(), "%s", + pTestCase->_correctBuffer[testId].c_str()); + // Exponenent representation + while ((exp = strstr(staging.data(), "E+")) != NULL + || (exp = strstr(staging.data(), "e+")) != NULL + || (exp = strstr(staging.data(), "E-")) != NULL + || (exp = strstr(staging.data(), "e-")) != NULL) { char correctExp[3]={0}; strncpy(correctExp,exp,2); - char* eCorrectBuffer = strstr((char*)pTestCase->_correctBuffer[testId].c_str(),correctExp); + // check if leading data is equal + int ret = strncmp(staging_correct.data(), staging.data(), + exp - staging.data()); + if (ret) return ret; + + char* eCorrectBuffer = strstr(staging_correct.data(), correctExp); if(eCorrectBuffer == NULL) return 1; @@ -1156,7 +1188,21 @@ size_t verifyOutputBuffer(char *analysisBuffer,testCase* pTestCase,size_t testId ++exp; while(*eCorrectBuffer == '0') ++eCorrectBuffer; - return strcmp(eCorrectBuffer,exp); + + copy_str = std::string(eCorrectBuffer); + std::snprintf(staging_correct.data(), staging_correct.size(), "%s", + copy_str.c_str()); + + copy_str = std::string(exp); + std::snprintf(staging.data(), staging.size(), "%s", copy_str.c_str()); + + if (strstr(staging.data(), "E+") != NULL + || strstr(staging.data(), "e+") != NULL + || strstr(staging.data(), "E-") != NULL + || strstr(staging.data(), "e-") != NULL) + continue; + + return strcmp(staging_correct.data(), copy_str.c_str()); } if (pTestCase->_correctBuffer[testId] == "inf") From 2be4ca5985aa0b093e85ef2e11730fbf6a2caac6 Mon Sep 17 00:00:00 2001 From: Marcin Hajder Date: Thu, 20 Jun 2024 10:17:11 +0200 Subject: [PATCH 2/5] Correction related to windows build in CI check --- test_conformance/printf/test_printf.cpp | 4 +++- 1 file changed, 3 insertions(+), 1 deletion(-) diff --git a/test_conformance/printf/test_printf.cpp b/test_conformance/printf/test_printf.cpp index 0fe517e14..0ce19bf61 100644 --- a/test_conformance/printf/test_printf.cpp +++ b/test_conformance/printf/test_printf.cpp @@ -18,10 +18,12 @@ #include "harness/stringHelpers.h" #include +#include #include -#include #include #include +#include +#include #if ! defined( _WIN32) #if defined(__APPLE__) From 2c9c11fe5499bf40e6889668e11019ff1e0e5909 Mon Sep 17 00:00:00 2001 From: Marcin Hajder Date: Fri, 5 Jul 2024 14:02:00 +0200 Subject: [PATCH 3/5] Corrections due to code review --- test_conformance/printf/test_printf.cpp | 67 +++++++++++++------------ test_conformance/printf/test_printf.h | 2 +- test_conformance/printf/util_printf.cpp | 6 ++- 3 files changed, 41 insertions(+), 34 deletions(-) diff --git a/test_conformance/printf/test_printf.cpp b/test_conformance/printf/test_printf.cpp index 0ce19bf61..334c07744 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 "harness/stringHelpers.h" +#include "harness/conversions.h" #include #include @@ -219,28 +220,25 @@ int waitForEvent(cl_event* event) return CL_SUCCESS; } -//----------------------------------------- -// Static helper functions definition -//----------------------------------------- - -float get_random_float(float low, float high, MTdata d) -{ - float t = (float)((double)genrand_int32(d) / (double)0xFFFFFFFF); - return (1.0f - t) * low + t * high; -} - //----------------------------------------- // makeMixedFormatPrintfProgram +// Generates in-flight printf kernel with format string including: +// -data before conversion flags (randomly generated ascii string) +// -randomly generated conversion flags (integer of floating point) +// -data after conversion flags (randomly generated ascii string). +// Moreover it generates suitable arguments. +// example: printf("zH, %u, %a, D+{gy\n", -929240879, 24295.671875f) //----------------------------------------- cl_program makeMixedFormatPrintfProgram(cl_kernel* kernel_ptr, const cl_context context, const cl_device_id device, const unsigned int testId, + const unsigned int testNum, const std::string& testname) { auto gen_char = [&]() { static const char dict[] = { - " !#$&()*+,-./" + " !#$&()*+,-./" "123456789:;<=>?@ABCDEFGHIJKLMNOPQRSTUVWXYZ[]^_`" "abcdefghijklmnopqrstuvwxyz{|}~" }; @@ -249,12 +247,12 @@ cl_program makeMixedFormatPrintfProgram(cl_kernel* kernel_ptr, std::array, 2> formats = { { { "%f", "%e", "%g", "%a", "%F", "%E", "%G", "%A" }, - { "%d", "%i", "%u", "%x" } } + { "%d", "%i", "%u", "%x", "%o" } } }; - std::vector data_befor(2 + genrand_int32(gMTdata) % 8); + std::vector data_before(2 + genrand_int32(gMTdata) % 8); std::vector data_after(2 + genrand_int32(gMTdata) % 8); - std::generate(data_befor.begin(), data_befor.end(), gen_char); + std::generate(data_before.begin(), data_before.end(), gen_char); std::generate(data_after.begin(), data_after.end(), gen_char); cl_uint num_args = 2 + genrand_int32(gMTdata) % 4; @@ -288,7 +286,7 @@ cl_program makeMixedFormatPrintfProgram(cl_kernel* kernel_ptr, << "(void)\n" "{\n" " printf(\""; - for (auto it : data_befor) + for (auto it : data_before) { format_str << it; ref_str << it; @@ -299,16 +297,16 @@ cl_program makeMixedFormatPrintfProgram(cl_kernel* kernel_ptr, for (cl_uint i = 0; i < num_args; i++) { - std::uint8_t type_ind = genrand_int32(gMTdata) % 2; + std::uint8_t is_int = genrand_int32(gMTdata) % 2; // Set CPU rounding mode to match that of the device - set_round(deviceRound, type_ind != 0 ? kint : kfloat); + set_round(deviceRound, is_int != 0 ? kint : kfloat); - std::string format = formats[type_ind][genrand_int32(gMTdata) - % formats[type_ind].size()]; + std::string format = + formats[is_int][genrand_int32(gMTdata) % formats[is_int].size()]; format_str << format << ", "; - if (type_ind) + if (is_int) { int arg = genrand_int32(gMTdata); args_str << str_sprintf("%d", arg) << ", "; @@ -325,14 +323,25 @@ cl_program makeMixedFormatPrintfProgram(cl_kernel* kernel_ptr, // Restore the original CPU rounding mode set_round(hostRound, kfloat); - args_str.seekp(-2, std::ios_base::end); - args_str << ");\n}\n"; - for (auto it : data_after) { format_str << it; ref_str << it; } + + { + std::ostringstream args_cpy; + args_cpy << args_str.str(); + args_cpy.seekp(-2, std::ios_base::end); + args_cpy << ")\n"; + log_info("%d) testing printf(\"%s\\n\", %s", testNum, + format_str.str().c_str(), args_cpy.str().c_str()); + } + + args_str.seekp(-2, std::ios_base::end); + args_str << ");\n}\n"; + + source_gen << format_str.str() << "\\n\"" << ", " << args_str.str(); @@ -472,10 +481,10 @@ cl_program makePrintfProgram(cl_kernel* kernel_ptr, const cl_context context, concat_kernel(sourceAddrSpace, sizeof(sourceAddrSpace) / sizeof(sourceAddrSpace[0])); } - else if (allTestCase[testId]->_type == TYPE_MIXED_FORMAT) + else if (allTestCase[testId]->_type == TYPE_MIXED_FORMAT_RANDOM) { return makeMixedFormatPrintfProgram(kernel_ptr, context, device, testId, - testname); + testNum, testname); } else { @@ -653,11 +662,7 @@ void logTestType(const unsigned testId, const unsigned testNum, .addrSpaceParameter); } } - else if (allTestCase[testId]->_type == TYPE_MIXED_FORMAT) - { - log_info("%d) testing printf with mixed format string\n", testNum); - } - else + else if (allTestCase[testId]->_type != TYPE_MIXED_FORMAT_RANDOM) { log_info("%d)testing printf(\"%s\"", testNum, allTestCase[testId] @@ -980,7 +985,7 @@ int test_address_space(cl_device_id deviceID, cl_context context, int test_mixed_format(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) { - return doTest(gQueue, gContext, TYPE_MIXED_FORMAT, deviceID); + return doTest(gQueue, gContext, TYPE_MIXED_FORMAT_RANDOM, deviceID); } int test_buffer_size(cl_device_id deviceID, cl_context context, diff --git a/test_conformance/printf/test_printf.h b/test_conformance/printf/test_printf.h index f2d667172..a2cd9ed2b 100644 --- a/test_conformance/printf/test_printf.h +++ b/test_conformance/printf/test_printf.h @@ -58,7 +58,7 @@ enum PrintfTestType TYPE_FORMAT_STRING, TYPE_VECTOR, TYPE_ADDRESS_SPACE, - TYPE_MIXED_FORMAT, + TYPE_MIXED_FORMAT_RANDOM, TYPE_COUNT }; diff --git a/test_conformance/printf/util_printf.cpp b/test_conformance/printf/util_printf.cpp index f3e2a24f0..6c61ffeac 100644 --- a/test_conformance/printf/util_printf.cpp +++ b/test_conformance/printf/util_printf.cpp @@ -1101,7 +1101,8 @@ testCase testCaseAddrSpace = { //---------------------------------------------------------- // placeholders for mixed-args-data test //---------------------------------------------------------- -// empty slots specifies number of tests +// number of records with empty objects is the number of tests during which the +// format string and reference string are generated std::vector printMixedFormatGenParameters(64, { { "" } }); @@ -1110,7 +1111,8 @@ std::vector correctBufferMixedFormat; //---------------------------------------------------------- // Test case for mixed-args //---------------------------------------------------------- -testCase testCaseMixedFormat = { TYPE_MIXED_FORMAT, correctBufferMixedFormat, +testCase testCaseMixedFormat = { TYPE_MIXED_FORMAT_RANDOM, + correctBufferMixedFormat, printMixedFormatGenParameters, NULL }; //------------------------------------------------------------------------------- From 2ade8d643f5b18905572aebc42c58799eaefde52 Mon Sep 17 00:00:00 2001 From: Marcin Hajder Date: Mon, 8 Jul 2024 13:24:01 +0200 Subject: [PATCH 4/5] Corrections due to code review --- test_conformance/printf/test_printf.cpp | 10 +++++----- test_conformance/printf/util_printf.cpp | 7 +++---- 2 files changed, 8 insertions(+), 9 deletions(-) diff --git a/test_conformance/printf/test_printf.cpp b/test_conformance/printf/test_printf.cpp index 334c07744..a4af25de0 100644 --- a/test_conformance/printf/test_printf.cpp +++ b/test_conformance/printf/test_printf.cpp @@ -224,7 +224,7 @@ int waitForEvent(cl_event* event) // makeMixedFormatPrintfProgram // Generates in-flight printf kernel with format string including: // -data before conversion flags (randomly generated ascii string) -// -randomly generated conversion flags (integer of floating point) +// -randomly generated conversion flags (integer or floating point) // -data after conversion flags (randomly generated ascii string). // Moreover it generates suitable arguments. // example: printf("zH, %u, %a, D+{gy\n", -929240879, 24295.671875f) @@ -238,7 +238,7 @@ cl_program makeMixedFormatPrintfProgram(cl_kernel* kernel_ptr, { auto gen_char = [&]() { static const char dict[] = { - " !#$&()*+,-./" + " \t!#$&()*+,-./" "123456789:;<=>?@ABCDEFGHIJKLMNOPQRSTUVWXYZ[]^_`" "abcdefghijklmnopqrstuvwxyz{|}~" }; @@ -247,7 +247,7 @@ cl_program makeMixedFormatPrintfProgram(cl_kernel* kernel_ptr, std::array, 2> formats = { { { "%f", "%e", "%g", "%a", "%F", "%E", "%G", "%A" }, - { "%d", "%i", "%u", "%x", "%o" } } + { "%d", "%i", "%u", "%x", "%o", "%X" } } }; std::vector data_before(2 + genrand_int32(gMTdata) % 8); std::vector data_after(2 + genrand_int32(gMTdata) % 8); @@ -982,7 +982,7 @@ int test_address_space(cl_device_id deviceID, cl_context context, return doTest(gQueue, gContext, TYPE_ADDRESS_SPACE, deviceID); } -int test_mixed_format(cl_device_id deviceID, cl_context context, +int test_mixed_format_random(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) { return doTest(gQueue, gContext, TYPE_MIXED_FORMAT_RANDOM, deviceID); @@ -1018,7 +1018,7 @@ test_definition test_list[] = { ADD_TEST(float), ADD_TEST(float_limits), ADD_TEST(octal), ADD_TEST(unsigned), ADD_TEST(hexadecimal), ADD_TEST(char), ADD_TEST(string), ADD_TEST(format_string), ADD_TEST(vector), - ADD_TEST(address_space), ADD_TEST(buffer_size), ADD_TEST(mixed_format), + ADD_TEST(address_space), ADD_TEST(buffer_size), ADD_TEST(mixed_format_random), }; const int test_num = ARRAY_SIZE( test_list ); diff --git a/test_conformance/printf/util_printf.cpp b/test_conformance/printf/util_printf.cpp index 6c61ffeac..eb913bce6 100644 --- a/test_conformance/printf/util_printf.cpp +++ b/test_conformance/printf/util_printf.cpp @@ -1099,10 +1099,9 @@ testCase testCaseAddrSpace = { //========================================================= //---------------------------------------------------------- -// placeholders for mixed-args-data test -//---------------------------------------------------------- -// number of records with empty objects is the number of tests during which the -// format string and reference string are generated +// Container related to mixed format tests. +// Empty records for which the format string and reference string are generated at run time. +// The size of this vector specifies the number of random tests that will be run. std::vector printMixedFormatGenParameters(64, { { "" } }); From d9bc1c768982867ac3ba85867ce9ca0db17e0339 Mon Sep 17 00:00:00 2001 From: Marcin Hajder Date: Mon, 8 Jul 2024 13:29:40 +0200 Subject: [PATCH 5/5] clang format correction --- test_conformance/printf/test_printf.cpp | 22 ++++++++++++++++------ test_conformance/printf/util_printf.cpp | 5 +++-- 2 files changed, 19 insertions(+), 8 deletions(-) diff --git a/test_conformance/printf/test_printf.cpp b/test_conformance/printf/test_printf.cpp index a4af25de0..424b42d3f 100644 --- a/test_conformance/printf/test_printf.cpp +++ b/test_conformance/printf/test_printf.cpp @@ -983,7 +983,7 @@ int test_address_space(cl_device_id deviceID, cl_context context, } int test_mixed_format_random(cl_device_id deviceID, cl_context context, - cl_command_queue queue, int num_elements) + cl_command_queue queue, int num_elements) { return doTest(gQueue, gContext, TYPE_MIXED_FORMAT_RANDOM, deviceID); } @@ -1014,11 +1014,21 @@ int test_buffer_size(cl_device_id deviceID, cl_context context, } test_definition test_list[] = { - ADD_TEST(int), ADD_TEST(half), ADD_TEST(half_limits), - ADD_TEST(float), ADD_TEST(float_limits), ADD_TEST(octal), - ADD_TEST(unsigned), ADD_TEST(hexadecimal), ADD_TEST(char), - ADD_TEST(string), ADD_TEST(format_string), ADD_TEST(vector), - ADD_TEST(address_space), ADD_TEST(buffer_size), ADD_TEST(mixed_format_random), + ADD_TEST(int), + ADD_TEST(half), + ADD_TEST(half_limits), + ADD_TEST(float), + ADD_TEST(float_limits), + ADD_TEST(octal), + ADD_TEST(unsigned), + ADD_TEST(hexadecimal), + ADD_TEST(char), + ADD_TEST(string), + ADD_TEST(format_string), + ADD_TEST(vector), + ADD_TEST(address_space), + ADD_TEST(buffer_size), + ADD_TEST(mixed_format_random), }; const int test_num = ARRAY_SIZE( test_list ); diff --git a/test_conformance/printf/util_printf.cpp b/test_conformance/printf/util_printf.cpp index eb913bce6..6e44b43fd 100644 --- a/test_conformance/printf/util_printf.cpp +++ b/test_conformance/printf/util_printf.cpp @@ -1100,8 +1100,9 @@ testCase testCaseAddrSpace = { //---------------------------------------------------------- // Container related to mixed format tests. -// Empty records for which the format string and reference string are generated at run time. -// The size of this vector specifies the number of random tests that will be run. +// Empty records for which the format string and reference string are generated +// at run time. The size of this vector specifies the number of random tests +// that will be run. std::vector printMixedFormatGenParameters(64, { { "" } });