| // |
| // Copyright (c) 2017 The Khronos Group Inc. |
| // |
| // Licensed under the Apache License, Version 2.0 (the "License"); |
| // you may not use this file except in compliance with the License. |
| // You may obtain a copy of the License at |
| // |
| // http://www.apache.org/licenses/LICENSE-2.0 |
| // |
| // Unless required by applicable law or agreed to in writing, software |
| // distributed under the License is distributed on an "AS IS" BASIS, |
| // WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. |
| // See the License for the specific language governing permissions and |
| // limitations under the License. |
| // |
| #include "harness/os_helpers.h" |
| #include "harness/typeWrappers.h" |
| #include "harness/stringHelpers.h" |
| #include "harness/conversions.h" |
| |
| #include <algorithm> |
| #include <array> |
| #include <cstdarg> |
| #include <cstdint> |
| #include <errno.h> |
| #include <memory> |
| #include <string.h> |
| #include <vector> |
| |
| #if ! defined( _WIN32) |
| #if defined(__APPLE__) |
| #include <sys/sysctl.h> |
| #endif |
| #include <unistd.h> |
| #define streamDup(fd1) dup(fd1) |
| #define streamDup2(fd1,fd2) dup2(fd1,fd2) |
| #endif |
| #include <limits.h> |
| #include <time.h> |
| #include "test_printf.h" |
| |
| #if defined(_WIN32) |
| #include <io.h> |
| #define streamDup(fd1) _dup(fd1) |
| #define streamDup2(fd1,fd2) _dup2(fd1,fd2) |
| #endif |
| |
| #include "harness/testHarness.h" |
| #include "harness/errorHelpers.h" |
| #include "harness/kernelHelpers.h" |
| #include "harness/parseParameters.h" |
| #include "harness/rounding_mode.h" |
| |
| #include <CL/cl_ext.h> |
| |
| typedef unsigned int uint32_t; |
| |
| |
| test_status InitCL( cl_device_id device ); |
| |
| namespace { |
| |
| //----------------------------------------- |
| // helper functions declaration |
| //----------------------------------------- |
| |
| //Stream helper functions |
| |
| //Associate stdout stream with the file(gFileName):i.e redirect stdout stream to the specific files (gFileName) |
| int acquireOutputStream(int* error); |
| |
| //Close the file(gFileName) associated with the stdout stream and disassociates it. |
| void releaseOutputStream(int fd); |
| |
| //Get analysis buffer to verify the correctess of printed data |
| void getAnalysisBuffer(char* analysisBuffer); |
| |
| //Kernel builder helper functions |
| |
| //Check if the test case is for kernel that has argument |
| int isKernelArgument(testCase* pTestCase, size_t testId); |
| |
| //Check if the test case treats %p format for void* |
| int isKernelPFormat(testCase* pTestCase, size_t testId); |
| |
| //----------------------------------------- |
| // Static functions declarations |
| //----------------------------------------- |
| // Make a program that uses printf for the given type/format, |
| 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 |
| int doTest(cl_command_queue queue, cl_context context, |
| const unsigned int testId, cl_device_id device); |
| |
| // Check if device supports long |
| bool isLongSupported(cl_device_id device_id); |
| |
| // Check if device address space is 64 bits |
| bool is64bAddressSpace(cl_device_id device_id); |
| |
| //Wait until event status is CL_COMPLETE |
| int waitForEvent(cl_event* event); |
| |
| //----------------------------------------- |
| // Definitions and initializations |
| //----------------------------------------- |
| |
| // Tests are broken into the major test which is based on the |
| // src and cmp type and their corresponding vector types and |
| // sub tests which is for each individual test. The following |
| // tracks the subtests |
| 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]; |
| |
| MTdataHolder gMTdata; |
| |
| // For the sake of proper logging of negative results |
| std::string gLatestKernelSource; |
| |
| //----------------------------------------- |
| // helper functions definition |
| //----------------------------------------- |
| |
| //----------------------------------------- |
| // acquireOutputStream |
| //----------------------------------------- |
| int acquireOutputStream(int* error) |
| { |
| int fd = streamDup(fileno(stdout)); |
| *error = 0; |
| if (!freopen(gFileName, "w", stdout)) |
| { |
| releaseOutputStream(fd); |
| *error = -1; |
| } |
| return fd; |
| } |
| |
| //----------------------------------------- |
| // releaseOutputStream |
| //----------------------------------------- |
| void releaseOutputStream(int fd) |
| { |
| fflush(stdout); |
| streamDup2(fd,fileno(stdout)); |
| close(fd); |
| } |
| |
| //----------------------------------------- |
| // printfCallBack |
| //----------------------------------------- |
| void CL_CALLBACK printfCallBack(const char* printf_data, size_t len, |
| size_t final, void* user_data) |
| { |
| fwrite(printf_data, 1, len, stdout); |
| } |
| |
| //----------------------------------------- |
| // getAnalysisBuffer |
| //----------------------------------------- |
| void getAnalysisBuffer(char* analysisBuffer) |
| { |
| FILE *fp; |
| memset(analysisBuffer,0,ANALYSIS_BUFFER_SIZE); |
| |
| fp = fopen(gFileName, "r"); |
| if (NULL == fp) |
| log_error("Failed to open analysis buffer ('%s')\n", strerror(errno)); |
| else if (0 |
| == std::fread(analysisBuffer, sizeof(analysisBuffer[0]), |
| ANALYSIS_BUFFER_SIZE, fp)) |
| log_error("No data read from analysis buffer\n"); |
| |
| fclose(fp); |
| } |
| |
| //----------------------------------------- |
| // isKernelArgument |
| //----------------------------------------- |
| int isKernelArgument(testCase* pTestCase, size_t testId) |
| { |
| return strcmp(pTestCase->_genParameters[testId].addrSpaceArgumentTypeQualifier,""); |
| } |
| //----------------------------------------- |
| // isKernelPFormat |
| //----------------------------------------- |
| int isKernelPFormat(testCase* pTestCase, size_t testId) |
| { |
| return strcmp(pTestCase->_genParameters[testId].addrSpacePAdd,""); |
| } |
| |
| //----------------------------------------- |
| // waitForEvent |
| //----------------------------------------- |
| int waitForEvent(cl_event* event) |
| { |
| cl_int status = clWaitForEvents(1, event); |
| if(status != CL_SUCCESS) |
| { |
| log_error("clWaitForEvents failed"); |
| return status; |
| } |
| |
| status = clReleaseEvent(*event); |
| if(status != CL_SUCCESS) |
| { |
| log_error("clReleaseEvent failed. (*event)"); |
| return status; |
| } |
| return CL_SUCCESS; |
| } |
| |
| //----------------------------------------- |
| // makeMixedFormatPrintfProgram |
| // Generates in-flight printf kernel with format string including: |
| // -data before conversion flags (randomly generated ascii string) |
| // -randomly generated conversion flags (integer or floating point) |
| // -data after conversion flags (randomly generated ascii string). |
| // Moreover it generates suitable arguments. |
| // example: printf("zH, %u, %a, D+{gy\n", -929240879, 24295.671875f) |
| //----------------------------------------- |
| cl_program makeMixedFormatPrintfProgram(cl_kernel* kernel_ptr, |
| const cl_context context, |
| const cl_device_id device, |
| const unsigned int testId, |
| const unsigned int testNum, |
| const std::string& testname) |
| { |
| auto gen_char = [&]() { |
| static const char dict[] = { |
| " \t!#$&()*+,-./" |
| "123456789:;<=>?@ABCDEFGHIJKLMNOPQRSTUVWXYZ[]^_`" |
| "abcdefghijklmnopqrstuvwxyz{|}~" |
| }; |
| return dict[genrand_int32(gMTdata) % ((int)sizeof(dict) - 1)]; |
| }; |
| |
| std::array<std::vector<std::string>, 2> formats = { |
| { { "%f", "%e", "%g", "%.13a", "%F", "%E", "%G", "%.13A" }, |
| { "%d", "%i", "%u", "%x", "%o", "%X" } } |
| }; |
| std::vector<char> data_before(2 + genrand_int32(gMTdata) % 8); |
| std::vector<char> data_after(2 + genrand_int32(gMTdata) % 8); |
| |
| std::generate(data_before.begin(), data_before.end(), gen_char); |
| std::generate(data_after.begin(), data_after.end(), gen_char); |
| |
| cl_uint num_args = 2 + genrand_int32(gMTdata) % 4; |
| |
| // Map device rounding to CTS rounding type |
| // get_default_rounding_mode supports RNE and RTZ |
| auto get_rounding = [](const cl_device_fp_config& fpConfig) { |
| if (fpConfig == CL_FP_ROUND_TO_NEAREST) |
| { |
| return kRoundToNearestEven; |
| } |
| else if (fpConfig == CL_FP_ROUND_TO_ZERO) |
| { |
| return kRoundTowardZero; |
| } |
| else |
| { |
| assert(false && "Unreachable"); |
| } |
| return kDefaultRoundingMode; |
| }; |
| |
| const RoundingMode hostRound = get_round(); |
| RoundingMode deviceRound = get_rounding(get_default_rounding_mode(device)); |
| |
| std::ostringstream format_str; |
| std::ostringstream ref_str; |
| std::ostringstream source_gen; |
| std::ostringstream args_str; |
| source_gen << "__kernel void " << testname |
| << "(void)\n" |
| "{\n" |
| " printf(\""; |
| for (auto it : data_before) |
| { |
| format_str << it; |
| ref_str << it; |
| } |
| format_str << ", "; |
| ref_str << ", "; |
| |
| |
| for (cl_uint i = 0; i < num_args; i++) |
| { |
| std::uint8_t is_int = genrand_int32(gMTdata) % 2; |
| |
| // Set CPU rounding mode to match that of the device |
| set_round(deviceRound, is_int != 0 ? kint : kfloat); |
| |
| std::string format = |
| formats[is_int][genrand_int32(gMTdata) % formats[is_int].size()]; |
| format_str << format << ", "; |
| |
| if (is_int) |
| { |
| int arg = genrand_int32(gMTdata); |
| args_str << str_sprintf("%d", arg) << ", "; |
| ref_str << str_sprintf(format, arg) << ", "; |
| } |
| else |
| { |
| const float max_range = 100000.f; |
| float arg = get_random_float(-max_range, max_range, gMTdata); |
| std::string arg_str = str_sprintf("%f", arg); |
| args_str << arg_str << "f, "; |
| float arg_deviceRound = std::stof(arg_str); |
| ref_str << str_sprintf(format, arg_deviceRound) << ", "; |
| } |
| } |
| // Restore the original CPU rounding mode |
| set_round(hostRound, kfloat); |
| |
| for (auto it : data_after) |
| { |
| format_str << it; |
| ref_str << it; |
| } |
| |
| { |
| std::ostringstream args_cpy; |
| args_cpy << args_str.str(); |
| args_cpy.seekp(-2, std::ios_base::end); |
| args_cpy << ")\n"; |
| log_info("%d) testing printf(\"%s\\n\", %s", testNum, |
| format_str.str().c_str(), args_cpy.str().c_str()); |
| } |
| |
| args_str.seekp(-2, std::ios_base::end); |
| args_str << ");\n}\n"; |
| |
| |
| source_gen << format_str.str() << "\\n\"" |
| << ", " << args_str.str(); |
| |
| std::string kernel_source = source_gen.str(); |
| const char* ptr = kernel_source.c_str(); |
| |
| cl_program program; |
| cl_int err = create_single_kernel_helper(context, &program, kernel_ptr, 1, |
| &ptr, testname.c_str()); |
| |
| gLatestKernelSource = kernel_source.c_str(); |
| |
| // Save the reference result |
| allTestCase[testId]->_correctBuffer.push_back(ref_str.str()); |
| |
| if (!program || err) |
| { |
| log_error("create_single_kernel_helper failed\n"); |
| return NULL; |
| } |
| |
| return program; |
| } |
| |
| //----------------------------------------- |
| // makePrintfProgram |
| //----------------------------------------- |
| 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; |
| char testname[256] = {0}; |
| char addrSpaceArgument[256] = {0}; |
| char addrSpacePAddArgument[256] = {0}; |
| char extension[128] = { 0 }; |
| |
| //Update testname |
| std::snprintf(testname, sizeof(testname), "%s%d", "test", testId); |
| |
| if (allTestCase[testId]->_type == TYPE_HALF |
| || allTestCase[testId]->_type == TYPE_HALF_LIMITS) |
| strcpy(extension, "#pragma OPENCL EXTENSION cl_khr_fp16 : enable\n"); |
| |
| |
| //Update addrSpaceArgument and addrSpacePAddArgument types, based on FULL_PROFILE/EMBEDDED_PROFILE |
| if(allTestCase[testId]->_type == TYPE_ADDRESS_SPACE) |
| { |
| std::snprintf(addrSpaceArgument, sizeof(addrSpaceArgument), "%s", |
| allTestCase[testId] |
| ->_genParameters[testNum] |
| .addrSpaceArgumentTypeQualifier); |
| |
| std::snprintf( |
| addrSpacePAddArgument, sizeof(addrSpacePAddArgument), "%s", |
| allTestCase[testId]->_genParameters[testNum].addrSpacePAdd); |
| } |
| |
| if (strlen(addrSpaceArgument) == 0) |
| std::snprintf(addrSpaceArgument, sizeof(addrSpaceArgument), "void"); |
| |
| // create program based on its type |
| |
| if(allTestCase[testId]->_type == TYPE_VECTOR) |
| { |
| if (strcmp(allTestCase[testId]->_genParameters[testNum].dataType, |
| "half") |
| == 0) |
| 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); |
| |
| gLatestKernelSource = |
| concat_kernel(sourceVec, sizeof(sourceVec) / sizeof(sourceVec[0])); |
| } |
| 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]), |
| sourceAddrSpace, testname); |
| |
| gLatestKernelSource = |
| concat_kernel(sourceAddrSpace, |
| sizeof(sourceAddrSpace) / sizeof(sourceAddrSpace[0])); |
| } |
| else if (allTestCase[testId]->_type == TYPE_MIXED_FORMAT_RANDOM) |
| { |
| return makeMixedFormatPrintfProgram(kernel_ptr, context, device, testId, |
| testNum, testname); |
| } |
| else |
| { |
| // 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); |
| |
| gLatestKernelSource = kernel_source.c_str(); |
| } |
| |
| if (!program || err) { |
| log_error("create_single_kernel_helper failed\n"); |
| return NULL; |
| } |
| |
| return program; |
| } |
| |
| //----------------------------------------- |
| // isLongSupported |
| //----------------------------------------- |
| bool isLongSupported(cl_device_id device_id) |
| { |
| size_t tempSize = 0; |
| cl_int status; |
| bool extSupport = true; |
| |
| // Device profile |
| status = clGetDeviceInfo( |
| device_id, |
| CL_DEVICE_PROFILE, |
| 0, |
| NULL, |
| &tempSize); |
| |
| if(status != CL_SUCCESS) |
| { |
| log_error("*** clGetDeviceInfo FAILED ***\n\n"); |
| return false; |
| } |
| |
| std::unique_ptr<char[]> profileType(new char[tempSize]); |
| if(profileType == NULL) |
| { |
| log_error("Failed to allocate memory(profileType)"); |
| return false; |
| } |
| |
| status = clGetDeviceInfo( |
| device_id, |
| CL_DEVICE_PROFILE, |
| sizeof(char) * tempSize, |
| profileType.get(), |
| NULL); |
| |
| |
| if(!strcmp("EMBEDDED_PROFILE",profileType.get())) |
| { |
| extSupport = is_extension_available(device_id, "cles_khr_int64"); |
| } |
| return extSupport; |
| } |
| //----------------------------------------- |
| // is64bAddressSpace |
| //----------------------------------------- |
| bool is64bAddressSpace(cl_device_id device_id) |
| { |
| cl_int status; |
| cl_uint addrSpaceB; |
| |
| // Device profile |
| status = clGetDeviceInfo( |
| device_id, |
| CL_DEVICE_ADDRESS_BITS, |
| sizeof(cl_uint), |
| &addrSpaceB, |
| NULL); |
| if(status != CL_SUCCESS) |
| { |
| log_error("*** clGetDeviceInfo FAILED ***\n\n"); |
| return false; |
| } |
| if(addrSpaceB == 64) |
| return true; |
| else |
| return false; |
| } |
| |
| //----------------------------------------- |
| // subtest_fail |
| //----------------------------------------- |
| void subtest_fail(const char* msg, ...) |
| { |
| if (msg) |
| { |
| va_list argptr; |
| va_start(argptr, msg); |
| vfprintf(stderr, msg, argptr); |
| va_end(argptr); |
| } |
| ++s_test_fail; |
| ++s_test_cnt; |
| } |
| |
| //----------------------------------------- |
| // logTestType - printout test details |
| //----------------------------------------- |
| |
| void logTestType(const unsigned testId, const unsigned testNum, |
| unsigned formatNum) |
| { |
| if (allTestCase[testId]->_type == TYPE_VECTOR) |
| { |
| log_info( |
| "%d)testing printf(\"%sv%s%s\",%s)\n", testNum, |
| allTestCase[testId]->_genParameters[testNum].vectorFormatFlag, |
| allTestCase[testId]->_genParameters[testNum].vectorSize, |
| allTestCase[testId]->_genParameters[testNum].vectorFormatSpecifier, |
| allTestCase[testId]->_genParameters[testNum].dataRepresentation); |
| } |
| else if (allTestCase[testId]->_type == TYPE_ADDRESS_SPACE) |
| { |
| if (isKernelArgument(allTestCase[testId], testNum)) |
| { |
| log_info("%d)testing kernel //argument %s \n printf(%s,%s)\n", |
| testNum, |
| allTestCase[testId] |
| ->_genParameters[testNum] |
| .addrSpaceArgumentTypeQualifier, |
| allTestCase[testId] |
| ->_genParameters[testNum] |
| .genericFormats[formatNum] |
| .c_str(), |
| allTestCase[testId] |
| ->_genParameters[testNum] |
| .addrSpaceParameter); |
| } |
| else |
| { |
| log_info("%d)testing kernel //variable %s \n printf(%s,%s)\n", |
| testNum, |
| allTestCase[testId] |
| ->_genParameters[testNum] |
| .addrSpaceVariableTypeQualifier, |
| allTestCase[testId] |
| ->_genParameters[testNum] |
| .genericFormats[formatNum] |
| .c_str(), |
| allTestCase[testId] |
| ->_genParameters[testNum] |
| .addrSpaceParameter); |
| } |
| } |
| else if (allTestCase[testId]->_type != TYPE_MIXED_FORMAT_RANDOM) |
| { |
| log_info("%d)testing printf(\"%s\"", testNum, |
| allTestCase[testId] |
| ->_genParameters[testNum] |
| .genericFormats[formatNum] |
| .c_str()); |
| if (allTestCase[testId]->_genParameters[testNum].dataRepresentation) |
| log_info(",%s", |
| allTestCase[testId] |
| ->_genParameters[testNum] |
| .dataRepresentation); |
| log_info(")\n"); |
| } |
| |
| fflush(stdout); |
| } |
| |
| //----------------------------------------- |
| // doTest |
| //----------------------------------------- |
| int doTest(cl_command_queue queue, cl_context context, |
| const unsigned int testId, cl_device_id device) |
| { |
| int err = TEST_FAIL; |
| |
| 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_LONG) && !isLongSupported(device)) |
| { |
| log_info("Skipping long because long is not supported.\n"); |
| return TEST_SKIPPED_ITSELF; |
| } |
| |
| if ((allTestCase[testId]->_type == TYPE_DOUBLE |
| || allTestCase[testId]->_type == TYPE_DOUBLE_LIMITS) |
| && !is_extension_available(device, "cl_khr_fp64")) |
| { |
| log_info("Skipping double because cl_khr_fp64 extension is not " |
| "supported.\n"); |
| return TEST_SKIPPED_ITSELF; |
| } |
| |
| auto& genParams = allTestCase[testId]->_genParameters; |
| |
| 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 (allTestCase[testId]->_type == TYPE_VECTOR) |
| { |
| auto is_vector_type_supported = [&](const char* type_name, |
| const char* ext_name) { |
| if ((strcmp(genParams[testNum].dataType, type_name) == 0) |
| && !is_extension_available(device, ext_name)) |
| { |
| log_info("Skipping %s because %s extension " |
| "is not supported.\n", |
| type_name, ext_name); |
| |
| s_test_skip++; |
| s_test_cnt++; |
| return false; |
| } |
| return true; |
| }; |
| |
| if (!is_vector_type_supported("half", "cl_khr_fp16")) continue; |
| |
| if (!is_vector_type_supported("double", "cl_khr_fp64")) continue; |
| |
| // 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; |
| } |
| } |
| |
| auto genParamsVec = allTestCase[testId]->_genParameters; |
| auto genFormatVec = genParamsVec[testNum].genericFormats; |
| |
| for (unsigned formatNum = 0; formatNum < genFormatVec.size(); |
| formatNum++) |
| { |
| logTestType(testId, testNum, formatNum); |
| |
| 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, device, testId, |
| testNum, formatNum); |
| 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)) |
| { |
| 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; |
| } |
| } |
| } |
| |
| 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) |
| { |
| 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 |
| { |
| 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 with kernel: " |
| "\n%s\n expected: %s\n got: %s\n", |
| gLatestKernelSource.c_str(), |
| allTestCase[testId]->_correctBuffer[testNum].c_str(), |
| _analysisBuffer); |
| continue; |
| } |
| } |
| else // 64-bit address space |
| { |
| if (0 |
| != verifyOutputBuffer(_analysisBuffer, allTestCase[testId], |
| testNum, out64)) |
| { |
| 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; |
| } |
| } |
| } |
| ++s_test_cnt; |
| } |
| |
| // 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(cl_device_id deviceID, cl_context context, cl_command_queue queue, |
| int num_elements) |
| { |
| return doTest(gQueue, gContext, TYPE_INT, deviceID); |
| } |
| |
| int test_long(cl_device_id deviceID, cl_context context, cl_command_queue queue, |
| int num_elements) |
| { |
| return doTest(gQueue, gContext, TYPE_LONG, deviceID); |
| } |
| |
| int test_half(cl_device_id deviceID, cl_context context, cl_command_queue queue, |
| int num_elements) |
| { |
| return doTest(gQueue, gContext, TYPE_HALF, deviceID); |
| } |
| |
| 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, deviceID); |
| } |
| |
| int test_float(cl_device_id deviceID, cl_context context, |
| cl_command_queue queue, int num_elements) |
| { |
| return doTest(gQueue, gContext, TYPE_FLOAT, deviceID); |
| } |
| |
| 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, deviceID); |
| } |
| |
| int test_double(cl_device_id deviceID, cl_context context, |
| cl_command_queue queue, int num_elements) |
| { |
| return doTest(gQueue, gContext, TYPE_DOUBLE, deviceID); |
| } |
| |
| int test_double_limits(cl_device_id deviceID, cl_context context, |
| cl_command_queue queue, int num_elements) |
| { |
| return doTest(gQueue, gContext, TYPE_DOUBLE_LIMITS, deviceID); |
| } |
| |
| int test_octal(cl_device_id deviceID, cl_context context, |
| cl_command_queue queue, int num_elements) |
| { |
| return doTest(gQueue, gContext, TYPE_OCTAL, deviceID); |
| } |
| |
| int test_unsigned(cl_device_id deviceID, cl_context context, |
| cl_command_queue queue, int num_elements) |
| { |
| return doTest(gQueue, gContext, TYPE_UNSIGNED, deviceID); |
| } |
| |
| int test_hexadecimal(cl_device_id deviceID, cl_context context, |
| cl_command_queue queue, int num_elements) |
| { |
| return doTest(gQueue, gContext, TYPE_HEXADEC, deviceID); |
| } |
| |
| int test_char(cl_device_id deviceID, cl_context context, cl_command_queue queue, |
| int num_elements) |
| { |
| return doTest(gQueue, gContext, TYPE_CHAR, deviceID); |
| } |
| |
| int test_string(cl_device_id deviceID, cl_context context, |
| cl_command_queue queue, int num_elements) |
| { |
| return doTest(gQueue, gContext, TYPE_STRING, deviceID); |
| } |
| |
| int test_format_string(cl_device_id deviceID, cl_context context, |
| cl_command_queue queue, int num_elements) |
| { |
| return doTest(gQueue, gContext, TYPE_FORMAT_STRING, deviceID); |
| } |
| |
| int test_vector(cl_device_id deviceID, cl_context context, |
| cl_command_queue queue, int num_elements) |
| { |
| return doTest(gQueue, gContext, TYPE_VECTOR, deviceID); |
| } |
| |
| 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, deviceID); |
| } |
| |
| int test_mixed_format_random(cl_device_id deviceID, cl_context context, |
| cl_command_queue queue, int num_elements) |
| { |
| return doTest(gQueue, gContext, TYPE_MIXED_FORMAT_RANDOM, deviceID); |
| } |
| |
| int test_length_specifier(cl_device_id deviceID, cl_context context, |
| cl_command_queue queue, int num_elements) |
| { |
| return doTest(gQueue, gContext, TYPE_LENGTH_SPECIFIER, deviceID); |
| } |
| |
| int test_buffer_size(cl_device_id deviceID, cl_context context, |
| cl_command_queue queue, int num_elements) |
| { |
| size_t printf_buff_size = 0; |
| const size_t printf_buff_size_req = !gIsEmbedded ? (1024 * 1024UL) : 1024UL; |
| const size_t config_size = sizeof(printf_buff_size); |
| cl_int err = CL_SUCCESS; |
| |
| err = clGetDeviceInfo(deviceID, CL_DEVICE_PRINTF_BUFFER_SIZE, config_size, |
| &printf_buff_size, NULL); |
| if (err != CL_SUCCESS) |
| { |
| log_error("Unable to query CL_DEVICE_PRINTF_BUFFER_SIZE"); |
| return TEST_FAIL; |
| } |
| |
| if (printf_buff_size < printf_buff_size_req) |
| { |
| log_error("CL_DEVICE_PRINTF_BUFFER_SIZE does not meet requirements"); |
| return TEST_FAIL; |
| } |
| |
| return TEST_PASS; |
| } |
| |
| test_definition test_list[] = { |
| ADD_TEST(int), |
| ADD_TEST(long), |
| ADD_TEST(half), |
| ADD_TEST(half_limits), |
| ADD_TEST(float), |
| ADD_TEST(float_limits), |
| ADD_TEST(double), |
| ADD_TEST(double_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(length_specifier), |
| }; |
| |
| const int test_num = ARRAY_SIZE( test_list ); |
| |
| //----------------------------------------- |
| // printUsage |
| //----------------------------------------- |
| static void printUsage(void) |
| { |
| log_info("test_printf: <optional: testnames> \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 |
| //----------------------------------------- |
| int main(int argc, const char* argv[]) |
| { |
| argc = parseCustomParam(argc, argv); |
| if (argc == -1) |
| { |
| return -1; |
| } |
| |
| const char ** argList = (const char **)calloc( argc, sizeof( char*) ); |
| |
| if( NULL == argList ) |
| { |
| log_error( "Failed to allocate memory for argList array.\n" ); |
| return 1; |
| } |
| |
| argList[0] = argv[0]; |
| size_t argCount = 1; |
| |
| for (int i=1; i < argc; ++i) { |
| const char *arg = argv[i]; |
| if (arg == NULL) |
| break; |
| |
| if (arg[0] == '-') |
| { |
| arg++; |
| while(*arg != '\0') |
| { |
| switch(*arg) { |
| case 'h': |
| printUsage(); |
| return 0; |
| default: |
| log_error( " <-- unknown flag: %c (0x%2.2x)\n)", *arg, *arg ); |
| printUsage(); |
| return 0; |
| } |
| arg++; |
| } |
| } |
| else { |
| argList[argCount] = arg; |
| argCount++; |
| } |
| } |
| |
| char* pcTempFname = get_temp_filename(); |
| if (pcTempFname != nullptr) |
| { |
| strncpy(gFileName, pcTempFname, sizeof(gFileName) - 1); |
| gFileName[sizeof(gFileName) - 1] = '\0'; |
| } |
| |
| free(pcTempFname); |
| |
| if (strlen(gFileName) == 0) |
| { |
| log_error("get_temp_filename failed\n"); |
| return -1; |
| } |
| |
| gMTdata = MTdataHolder(gRandomSeed); |
| |
| int err = runTestHarnessWithCheck( argCount, argList, test_num, test_list, true, 0, InitCL ); |
| |
| if(gQueue) |
| { |
| int error = clFinish(gQueue); |
| if (error) { |
| log_error("clFinish failed: %d\n", error); |
| } |
| } |
| |
| if(clReleaseCommandQueue(gQueue)!=CL_SUCCESS) |
| log_error("clReleaseCommandQueue\n"); |
| if(clReleaseContext(gContext)!= CL_SUCCESS) |
| log_error("clReleaseContext\n"); |
| |
| |
| free(argList); |
| remove(gFileName); |
| return err; |
| } |
| |
| test_status InitCL( cl_device_id device ) |
| { |
| uint32_t device_frequency = 0; |
| uint32_t compute_devices = 0; |
| |
| int err; |
| gFd = acquireOutputStream(&err); |
| if (err != 0) |
| { |
| log_error("Error while redirection stdout to file"); |
| return TEST_FAIL; |
| } |
| |
| size_t config_size = sizeof( device_frequency ); |
| #if MULTITHREAD |
| if( (err = clGetDeviceInfo(device, CL_DEVICE_MAX_COMPUTE_UNITS, config_size, &compute_devices, NULL )) ) |
| #endif |
| compute_devices = 1; |
| |
| config_size = sizeof(device_frequency); |
| if((err = clGetDeviceInfo(device, CL_DEVICE_MAX_CLOCK_FREQUENCY, config_size, &device_frequency, NULL ))) |
| device_frequency = 1; |
| |
| releaseOutputStream(gFd); |
| |
| log_info( "\nCompute Device info:\n" ); |
| log_info( "\tProcessing with %d devices\n", compute_devices ); |
| log_info( "\tDevice Frequency: %d MHz\n", device_frequency ); |
| |
| printDeviceHeader( device ); |
| |
| PrintArch(); |
| |
| auto version = get_device_cl_version(device); |
| auto expected_min_version = Version(1, 2); |
| if (version < expected_min_version) |
| { |
| version_expected_info("Test", "OpenCL", |
| expected_min_version.to_string().c_str(), |
| version.to_string().c_str()); |
| return TEST_SKIP; |
| } |
| |
| gFd = acquireOutputStream(&err); |
| if (err != 0) |
| { |
| log_error("Error while redirection stdout to file"); |
| return TEST_FAIL; |
| } |
| cl_context_properties printf_properties[] = { |
| CL_PRINTF_CALLBACK_ARM, (cl_context_properties)printfCallBack, |
| CL_PRINTF_BUFFERSIZE_ARM, ANALYSIS_BUFFER_SIZE, 0 |
| }; |
| |
| cl_context_properties* props = NULL; |
| |
| if(is_extension_available(device, "cl_arm_printf")) |
| { |
| props = printf_properties; |
| } |
| |
| gContext = clCreateContext(props, 1, &device, notify_callback, NULL, NULL); |
| checkNull(gContext, "clCreateContext"); |
| |
| gQueue = clCreateCommandQueue(gContext, device, 0, NULL); |
| checkNull(gQueue, "clCreateCommandQueue"); |
| |
| releaseOutputStream(gFd); |
| |
| if (is_extension_available(device, "cl_khr_fp16")) |
| { |
| const cl_device_fp_config fpConfigHalf = |
| get_default_rounding_mode(device, CL_DEVICE_HALF_FP_CONFIG); |
| if (fpConfigHalf == CL_FP_ROUND_TO_NEAREST) |
| { |
| half_rounding_mode = CL_HALF_RTE; |
| } |
| else if (fpConfigHalf == CL_FP_ROUND_TO_ZERO) |
| { |
| half_rounding_mode = CL_HALF_RTZ; |
| } |
| else |
| { |
| log_error("Error while acquiring half rounding mode"); |
| } |
| } |
| |
| // Generate reference results |
| generateRef(device); |
| |
| return TEST_PASS; |
| } |
