diff --git a/test_conformance/printf/test_printf.cpp b/test_conformance/printf/test_printf.cpp index e43e302f1e..21b48efa45 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,16 @@ 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, const unsigned int testNum, + cl_device_id device); // Check if device supports long static bool isLongSupported(cl_device_id device_id); @@ -206,7 +213,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; @@ -223,7 +235,10 @@ static cl_program makePrintfProgram(cl_kernel *kernel_ptr, const cl_context cont "(void)\n", "{\n" " printf(\"", - allTestCase[testId]->_genParameters[testNum].genericFormat, + allTestCase[testId] + ->_genParameters[testNum] + .genericFormats[formatNum] + .c_str(), "\\n\",", allTestCase[testId]->_genParameters[testNum].dataRepresentation, ");", @@ -254,12 +269,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, "); ", @@ -399,22 +422,17 @@ static bool is64bAddressSpace(cl_device_id device_id) //----------------------------------------- // doTest //----------------------------------------- -static int doTest(cl_command_queue queue, cl_context context, const unsigned int testId, const unsigned int testNum, cl_device_id device) -{ - 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"); - return TEST_SKIPPED_ITSELF; - } - - if(allTestCase[testId]->_type == TYPE_VECTOR) +static int doTest(cl_command_queue queue, cl_context context, + const unsigned int testId, const unsigned int testNum, + cl_device_id device) +{ + int err = TEST_FAIL; + for (unsigned formatNum = 0; formatNum + < allTestCase[testId]->_genParameters[testNum].genericFormats.size(); + formatNum++) { - if ((strcmp(allTestCase[testId]->_genParameters[testNum].dataType, - "half") - == 0) + 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 " @@ -422,171 +440,231 @@ static int doTest(cl_command_queue queue, cl_context context, const unsigned int return TEST_SKIPPED_ITSELF; } - 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)) + 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].genericFormat,allTestCase[testId]->_genParameters[testNum].addrSpaceParameter); + 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; + } + + 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, - allTestCase[testId]->_genParameters[testNum].genericFormat,allTestCase[testId]->_genParameters[testNum].addrSpaceParameter); + log_info("%d)testing printf(\"%s\",%s)\n", testNum, + allTestCase[testId] + ->_genParameters[testNum] + .genericFormats[formatNum] + .c_str(), + allTestCase[testId] + ->_genParameters[testNum] + .dataRepresentation); } - } - 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" ); - return 0; - } + // 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; + } - // 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; + } - 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; - } + 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) + { + ++s_test_fail; + ++s_test_cnt; + return -1; + } - //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) { - log_error("clCreateBuffer failed\n"); - goto exit; + 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"); + continue; + } + err = clSetKernelArg(kernel, 0, sizeof(cl_mem), &d_a); + if (err != CL_SUCCESS) + { + log_error("clSetKernelArg failed\n"); + continue; + } } - err = clSetKernelArg(kernel, 0, sizeof(cl_mem), &d_a); - if(err!= CL_SUCCESS) { - log_error("clSetKernelArg failed\n"); - goto exit; + // 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"); + continue; + } + err = clSetKernelArg(kernel, 1, sizeof(cl_mem), &d_out); + if (err != CL_SUCCESS) + { + log_error("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) { - log_error("clCreateBuffer failed\n"); - goto exit; - } - err = clSetKernelArg(kernel, 1, sizeof(cl_mem), &d_out); - if(err!= CL_SUCCESS) { - log_error("clSetKernelArg failed\n"); - goto exit; - } + log_error("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); + log_error("\n clEnqueueNDRangeKernel failed errcode:%d\n", err); + ++s_test_fail; + 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); + log_error("clFlush failed\n"); + 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); + if (err != CL_SUCCESS) + { + log_error("waitforEvent failed\n"); + 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 + // + // Get the output printed from the kernel to _analysisBuffer + // and verify its correctness + getAnalysisBuffer(_analysisBuffer); + if (!is64bAddressSpace(device)) // 32-bit address space { - clEnqueueReadBuffer(queue, d_out, CL_TRUE, 0, sizeof(cl_int),&out32, - 0, NULL, NULL); + if (0 + != verifyOutputBuffer(_analysisBuffer, allTestCase[testId], + testNum, (cl_ulong)out32)) + err = ++s_test_fail; } else //64-bit address space { - clEnqueueReadBuffer(queue, d_out, CL_TRUE, 0, sizeof(cl_ulong),&out64, - 0, NULL, NULL); + if (0 + != verifyOutputBuffer(_analysisBuffer, allTestCase[testId], + testNum, out64)) + err = ++s_test_fail; } } - // - //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; } diff --git a/test_conformance/printf/test_printf.h b/test_conformance/printf/test_printf.h index 8eb2a03249..d8e9c2a253 100644 --- a/test_conformance/printf/test_printf.h +++ b/test_conformance/printf/test_printf.h @@ -62,7 +62,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 6b310a994d..2257659d5e 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" }, }; //--------------------------------------------------------- @@ -203,16 +205,22 @@ testCase testCaseHalf = { std::vector printHalfLimitsGenParameters = { // Infinity (1.0/0.0) + { { "%f", "%e", "%g", "%a" }, "1.0h/0.0h" }, - { "%f", "1.0h/0.0h" }, + // NaN + { { "%f", "%e", "%g", "%a" }, "nan((ushort)0)" }, // NaN + { { "%f", "%e", "%g", "%a" }, "acospi(2.0h)" }, - { "%f", "sqrt(-1.0h)" }, + // Infinity (1.0/0.0) + { { "%F", "%E", "%G", "%A" }, "1.0h/0.0h" }, // NaN - { "%f", "acospi(2.0h)" } + { { "%F", "%E", "%G", "%A" }, "nan((ushort)0)" }, + // NaN + { { "%F", "%E", "%G", "%A" }, "acospi(2.0h)" } }; //-------------------------------------------------------- @@ -224,9 +232,15 @@ std::vector correctBufferHalfLimits = { "inf", - "-nan", + "nan", + + "nan", + + "INF", + + "NAN", - "nan" + "NAN" }; @@ -265,77 +279,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" }, }; //--------------------------------------------------------- @@ -376,16 +402,22 @@ testCase testCaseFloat = { std::vector printFloatLimitsGenParameters = { // Infinity (1.0/0.0) + { { "%f", "%e", "%g", "%a" }, "1.0f/0.0f" }, - { "%f", "1.0f/0.0f" }, + // NaN + { { "%f", "%e", "%g", "%a" }, "nan(0U)" }, // NaN + { { "%f", "%e", "%g", "%a" }, "acospi(2.0f)" }, - { "%f", "sqrt(-1.0f)" }, + // Infinity (1.0/0.0) + { { "%F", "%E", "%G", "%A" }, "1.0f/0.0f" }, // NaN - { "%f", "acospi(2.0f)" } + { { "%F", "%E", "%G", "%A" }, "nan(0U)" }, + // NaN + { { "%F", "%E", "%G", "%A" }, "acospi(2.0f)" } }; //-------------------------------------------------------- @@ -397,9 +429,15 @@ std::vector correctBufferFloatLimits = { "inf", - "-nan", + "nan", + + "nan", - "nan" + "INF", + + "NAN", + + "NAN" }; @@ -437,21 +475,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 +532,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 +580,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 +640,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" } }; @@ -674,15 +713,15 @@ 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" }, "\"%%\"" }, }; //--------------------------------------------------------- @@ -741,27 +780,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 +860,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" } }; @@ -955,17 +1038,25 @@ size_t verifyOutputBuffer(char *analysisBuffer,testCase* pTestCase,size_t testId ++eCorrectBuffer; return strcmp(eCorrectBuffer,exp); } - if(!strcmp(pTestCase->_correctBuffer[testId].c_str(),"inf")) - return strcmp(analysisBuffer,"inf")&&strcmp(analysisBuffer,"infinity")&&strcmp(analysisBuffer,"1.#INF00")&&strcmp(analysisBuffer,"Inf"); - if(!strcmp(pTestCase->_correctBuffer[testId].c_str(),"nan") || !strcmp(pTestCase->_correctBuffer[testId].c_str(),"-nan")) { - return strcmp(analysisBuffer,"nan")&&strcmp(analysisBuffer,"-nan")&&strcmp(analysisBuffer,"1.#IND00")&&strcmp(analysisBuffer,"-1.#IND00")&&strcmp(analysisBuffer,"NaN")&&strcmp(analysisBuffer,"nan(ind)")&&strcmp(analysisBuffer,"nan(snan)")&&strcmp(analysisBuffer,"-nan(ind)"); - } - return strcmp(analysisBuffer,pTestCase->_correctBuffer[testId].c_str()); + + if (pTestCase->_correctBuffer[testId] == "inf") + return strcmp(analysisBuffer, "inf") + && strcmp(analysisBuffer, "infinity"); + else if (pTestCase->_correctBuffer[testId] == "INF") + return strcmp(analysisBuffer, "INF") + && strcmp(analysisBuffer, "INFINITY"); + else if (pTestCase->_correctBuffer[testId] == "nan") + return strcmp(analysisBuffer, "nan") && strcmp(analysisBuffer, "-nan"); + else if (pTestCase->_correctBuffer[testId] == "NAN") + return strcmp(analysisBuffer, "NAN") && strcmp(analysisBuffer, "-NAN"); + + return strcmp(analysisBuffer, pTestCase->_correctBuffer[testId].c_str()); } 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 +1064,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); } /*