/*
 * oclh_h_ws_base.c
 *      Author: havock
 */
#include <string.h>
#include <oclh_h_base_defs.h>
#include <oclh_h_settings.h>
#include <oclh_h_base_dev_clapi_wrappers.h>
#include <oclh_h_ws_base.h>
#include <oclh_h_ws_base_log.h>
#include <oclh_h_ws_log_clapi_reps.h>
#include <oclh_h_clapi_callbacks.h>
#include <oclh_h_internals.h>
#include <oclh_h_externals.h>

_GHT_WRKSET _ghf_declWS(void) {
    _GHT_WRKSET wSet={
        .pwSetAddr=NULL,
        .Queue=NULL,         .Program=NULL,
        .szNofCmpUnits=0ul,  .szNofRPThrds=0ul,  .szMaxWGsz=0ul,
        .szNofWIinWG=0ul,    .szNofAllWI=0ul,
        .pszMemBytes=NULL,
        .Log=_ghf_declLog(), .i32EigenLogFlag=0,
        .APIErr=0,
        .pKernels=NULL,      .KerErr=0,          .d_pKerErr=NULL,
        .pvDat=NULL,         .pfnDatCleaner=NULL
    };
    return(wSet);
}

int32_t _ghf_genrWS(      _GHT_WRKSET* const  pwSet,
                    const _GHT_LOG            Log,
                    const int32_t             i32ExclusiveLogFlag,
                    const cl_device_id        clWrkDev,
                          char**       const  ppcCLProgramSources,
                    const _GHE_SRCTYPE        SourceType,
                    const char*        const  pcOCLBuildOpts,
                    const _GHE_LOGLVL         LogLvl,
                    const _GHE_BUILD_LOG_MODE BuildLogMode
#if defined(__OCLH_BUILDER_FLAG) || \
    defined(__OCLH_COMPILER_ONLY_FLAG) || \
    defined(__OCLH_LINKER_ONLY_FLAG)
                  ,       char*        const  pcOutputPrefix
#endif /* defined(__OCLH_BUILDER_FLAG) ||
          defined(__OCLH_COMPILER_ONLY_FLAG) ||
          defined(__OCLH_LINKER_ONLY_FLAG) */
                                                            ) {
/**/
#define __CLAPI_ERR_ROUTINE_WITH_RET(___CLAPI_CALL) \
{   const cl_int tmp=pwSet->APIErr; \
    char pcLogMsg[_GHM_MAXLEN_OF_LOGMSG]; \
    snprintf(pcLogMsg, _GHM_MAXLEN_OF_LOGMSG, \
             "%s/%s", __func__, ___CLAPI_CALL); \
    _ghf_logWS_APIErr(*pwSet, pcLogMsg); \
    _ghf_wipeWS(pwSet); \
    return(tmp); }
#define __CLAPI_CREATE_PROGRAM_ERR_TO_STDOUT(___CLAPI_CALL) \
{   FILE* pTmpFilePtr=pwSet->Log.pfOut; \
    pwSet->Log.pfOut=stdout; \
    _ghf_logWS_DevInfoShort(*pwSet,_ghf_getWS_Dev(*pwSet),NULL); \
    snprintf(pcLogMsg,_GHM_MAXLEN_OF_LOGMSG,___CLAPI_CALL); \
    _ghf_logWS_APIErr(*pwSet, pcLogMsg); _ghf_logWS_Delim(*pwSet); \
    pwSet->Log.pfOut=pTmpFilePtr; }
/**/
    char pcLogMsg[_GHM_MAXLEN_OF_LOGMSG];
    if(!pwSet) return(_GHM_NULL_POINTER_RECEIVED_ERROR);
    _ghf_wipeWS(pwSet);
    pwSet->pvDat=NULL;
    pwSet->pfnDatCleaner=NULL;
    pwSet->pwSetAddr=pwSet;
    pwSet->APIErr=0;
    if(i32ExclusiveLogFlag || !Log.pfOut || !Log.pMtx) {
        _GHT_LOG ExclLog=_ghf_declLog();
        char pcFileName[_GHM_MAX_PATH_LENGTH]="\0";
        int32_t err=0;
        {
            time_t timeRaw; struct tm* stTime; time(&timeRaw);
            stTime=localtime(&timeRaw);
            snprintf(pcFileName,_GHM_MAX_PATH_LENGTH,
                     "%s/%04d%02d%02d%02d%02d%02d-wset-%lx.log",
                     _GHM_LOG_PATH,
                     stTime->tm_year+1900,stTime->tm_mon+1,stTime->tm_mday,
                     stTime->tm_hour,stTime->tm_min,stTime->tm_sec,
                     (uint64_t) pwSet);
        }
        err=_ghf_genrLog(&ExclLog,pcFileName); if(err) return(err);
        pwSet->Log=ExclLog;
        pwSet->i32EigenLogFlag=1;
    } else { pwSet->Log=Log; pwSet->i32EigenLogFlag=0; }
    snprintf(pcLogMsg,_GHM_MAXLEN_OF_LOGMSG,
             "(Re)build of workset 0x%lx", (uint64_t) pwSet);
    _ghf_logWS_Hdr(*pwSet, pcLogMsg);
    {
        int32_t i32Zero=0;
        { /* build queue as a key object (assembly point) */
            cl_context clCtx=NULL;
            cl_platform_id clPlatf=
                (cl_platform_id)_ghf_getDevInf_uintptr(clWrkDev,
                                                       CL_DEVICE_PLATFORM,
                                                       &pwSet->APIErr);
            if(!clPlatf || pwSet->APIErr)
                __CLAPI_ERR_ROUTINE_WITH_RET("clGetDeviceInfo")
            {
                /*
                 * TODO:  NVidia GPU после создания контекста теряет настройку
                 * CL_CONTEXT_INTEROP_USER_SYNC, в связи с чем установить её
                 * состояние становится затруднительным. Возможно стоит
                 * сохранять её в полях workset'а
                 */
                cl_context_properties pclCtxProps[5]={
                    CL_CONTEXT_PLATFORM, (cl_context_properties)clPlatf,
                    CL_CONTEXT_INTEROP_USER_SYNC, CL_FALSE, 0l };
                clCtx=
                    clCreateContext(pclCtxProps,   1u,   &clWrkDev,
                                    _ghf_CtxEvent, pwSet,&pwSet->APIErr);
                    /* TODO: ^[Valg] 8 bytes in 1 blocks are definitely lost
                     * in loss record 12 of 1,583 by Intel library
                     * Intel::OpenCL::TaskExecutor::TBBTaskExecutor::Init */
                if(pwSet->APIErr) {
                    /* обработка ситуации, при которой не обеспечивается
                     * автоматическую синхронизация между разными API */
                    pclCtxProps[2]=(cl_context_properties)0l;
                    if(clCtx) pwSet->APIErr=clReleaseContext(clCtx);
                    clCtx=clCreateContext(pclCtxProps,   1u,   &clWrkDev,
                                          _ghf_CtxEvent, pwSet,&pwSet->APIErr);
                    /* TODO: ^[Valg] 4 bytes in 1 blocks are definitely lost
                     * in loss record 1 of 1,587 */
                    if(pwSet->APIErr)
                        __CLAPI_ERR_ROUTINE_WITH_RET("clCreateContext")
                }
            }
            _ghf_logWS_ContextInfo(*pwSet,clCtx,LogLvl);
            /* command queue */
            if(_GHM_DEBUG_N_PROFILE)
                pwSet->Queue=
                    clCreateCommandQueue(clCtx, clWrkDev,
                                         CL_QUEUE_PROFILING_ENABLE,
                                         &pwSet->APIErr);
            else
                pwSet->Queue=
                    clCreateCommandQueue(clCtx, clWrkDev, 0ul, &pwSet->APIErr);
            if(pwSet->APIErr)
                __CLAPI_ERR_ROUTINE_WITH_RET("clCreateCommandQueue")
            /* TODO: make queue info logging */
        }
        if(LogLvl==_GHE_LOG_DEFAULT)
            _ghf_logWS_DevInfoShort(*pwSet, _ghf_getWS_Dev(*pwSet), NULL);
        if(ppcCLProgramSources) { /* program */
            cl_device_id clWrkDev=_ghf_getWS_Dev(*pwSet);
            switch(SourceType) {
            case _GHE_HL_LISTINGS:
              { cl_uint cluNofListings=0u;
                while(((char**)ppcCLProgramSources)[cluNofListings]) {
                    if(_GHM_SHOW_PROGRAM_LISTING_IN_LOG) {
                        char pcLogMsg[64];
                        snprintf(pcLogMsg, 64ul,
                                 "Program: #%u", cluNofListings);
                        _ghf_logWS_Hdr(*pwSet, pcLogMsg);
                        _ghf_logWS_Raw(*pwSet,
                                 ((char**)ppcCLProgramSources)[cluNofListings]);
                        _ghf_logWS_Hdr(*pwSet, "End of program");
                    }
                    cluNofListings++;
                }
                pwSet->Program=
                    clCreateProgramWithSource(
                              _ghf_getWS_Ctx(*pwSet),cluNofListings,
                              (const char**)ppcCLProgramSources,
                              NULL,
                              /* TODO: стоило бы использовать размеры вместо
                               * NULL, но по неясной причине передача длин
                               * исходных текстов препятствует созданию
                               * ядерных функций. Необходимо разобраться в
                               * причинах. */
                          /*((size_t**)ppcCLProgramSources)[cluNofListings+1],*/
                              &pwSet->APIErr);
                if(pwSet->APIErr) {
#ifdef __OCLH_BUILD_LOG_TO_STDOUT_FLAG
              __CLAPI_CREATE_PROGRAM_ERR_TO_STDOUT("clCreateProgramWithSource");
#endif /* __OCLH_BUILD_LOG_TO_STDOUT_FLAG */
                    __CLAPI_ERR_ROUTINE_WITH_RET("clCreateProgramWithSource")
                }
              }
            break;
            case _GHE_IR_LISTINGS:
              { /* TODO: make IR output and processing */
                _ghf_logWS_Msg(*pwSet, "_GHE_IR_LISTINGS");
                _ghf_wipeWS(pwSet); return(1); 
              }
            break;
            case _GHE_SEPARATED_OBJECTS:
              { cl_device_id cldev=_ghf_getWS_Dev(*pwSet);
                cl_program* pPrograms=NULL;
                cl_uint cluNofBinaries=0u,
                        i=0u;
                while(((char**)ppcCLProgramSources)[cluNofBinaries])
                    cluNofBinaries++;
                pPrograms=malloc(sizeof(cl_program)*cluNofBinaries);
                if(!pPrograms) {
                    char pcLogMsg[_GHM_MAXLEN_OF_LOGMSG];
                    snprintf(pcLogMsg, _GHM_MAXLEN_OF_LOGMSG,
                             "%s/Host memory allocation failed", __func__);
                    _ghf_logWS_Msg(*pwSet, pcLogMsg); _ghf_wipeWS(pwSet);
                    return(_GHM_HOST_MEMALLOC_ERROR);
                }
                for(i=0u; i<cluNofBinaries; i++) pPrograms[i]=NULL;
                for(i=0u; i<cluNofBinaries; i++) {
                    cl_int cliBinStatus=CL_SUCCESS;
                    pPrograms[i]=
                        clCreateProgramWithBinary(
                          _ghf_getWS_Ctx(*pwSet),1u,&cldev,
                          &((size_t**)ppcCLProgramSources)[cluNofBinaries+1][i],
                          &((const uint8_t**)ppcCLProgramSources)[i],
                          &cliBinStatus, &pwSet->APIErr);
                    if(pwSet->APIErr) {
                        cl_int err=pwSet->APIErr;
                        for(i=0u; i<cluNofBinaries; i++)
                            if(pPrograms[i])
                                pwSet->APIErr=clReleaseProgram(pPrograms[i]);
                        _ghf_freeHostZ(&pPrograms);
                        pwSet->APIErr=err;
#ifdef __OCLH_BUILD_LOG_TO_STDOUT_FLAG
              __CLAPI_CREATE_PROGRAM_ERR_TO_STDOUT("clCreateProgramWithBinary");
#endif /* __OCLH_BUILD_LOG_TO_STDOUT_FLAG */
                       __CLAPI_ERR_ROUTINE_WITH_RET("clCreateProgramWithBinary")
                    }
                }
                pwSet->Program=clLinkProgram(_ghf_getWS_Ctx(*pwSet), 1u, &cldev,
                                             pcOCLBuildOpts?pcOCLBuildOpts:"",
                                             cluNofBinaries, pPrograms,
                                             _ghf_LinkEvent, pwSet,
                                             &pwSet->APIErr);
                if(pwSet->APIErr) {
                    cl_int err=pwSet->APIErr;
                    for(i=0u; i<cluNofBinaries; i++)
                        if(pPrograms[i])
                            pwSet->APIErr=clReleaseProgram(pPrograms[i]);
                    _ghf_freeHostZ(&pPrograms);
                    pwSet->APIErr=err;
#ifdef __OCLH_BUILD_LOG_TO_STDOUT_FLAG
                    __CLAPI_CREATE_PROGRAM_ERR_TO_STDOUT("clLinkProgram");
#endif /* __OCLH_BUILD_LOG_TO_STDOUT_FLAG */
                    __CLAPI_ERR_ROUTINE_WITH_RET("clLinkProgram")
                }
                for(i=0u; i<cluNofBinaries; i++)
                    if(pPrograms[i])
                        pwSet->APIErr=clReleaseProgram(pPrograms[i]);
                _ghf_freeHostZ(&pPrograms);
              }
            break;
            case _GHE_LINKED_OBJECTS:
              { cl_device_id cldev=_ghf_getWS_Dev(*pwSet);
                cl_uint cluNofBinaries=0u;
                cl_int cliBinStatus=CL_SUCCESS;
                while(((char**)ppcCLProgramSources)[cluNofBinaries])
                    cluNofBinaries++;
                pwSet->Program=
                    clCreateProgramWithBinary(
                              _ghf_getWS_Ctx(*pwSet),1u,&cldev,
                              ((size_t**)ppcCLProgramSources)[cluNofBinaries+1],
                              (const uint8_t**)ppcCLProgramSources,
                              &cliBinStatus, &pwSet->APIErr);
                if(pwSet->APIErr) {
#ifdef __OCLH_BUILD_LOG_TO_STDOUT_FLAG
              __CLAPI_CREATE_PROGRAM_ERR_TO_STDOUT("clCreateProgramWithBinary");
#endif /* __OCLH_BUILD_LOG_TO_STDOUT_FLAG */
                    __CLAPI_ERR_ROUTINE_WITH_RET("clCreateProgramWithBinary")
                }
              }
            break;
            default:
              { char pcLogMsg[_GHM_MAXLEN_OF_LOGMSG];
                snprintf(pcLogMsg, _GHM_MAXLEN_OF_LOGMSG,
                         "Unknown program source");
                _ghf_logWS_APIErr(*pwSet, pcLogMsg);
                _ghf_wipeWS(pwSet);
                return(_GHM_UNKNOWN_PROGRAM_SOURCE_ERROR);
              }
            }
            {
                cl_int cliBuildCLAPIErr=CL_SUCCESS;
                char *pcBldOpts=NULL;
                size_t szBOLen=
                    strlen(_GHM_OCLH_COMMON_BUILD_OPTIONS)+
                    strlen(pcOCLBuildOpts?pcOCLBuildOpts:"")+2ul;
                pcBldOpts=malloc(szBOLen);
                if(!pcBldOpts) {
                    char pcLogMsg[_GHM_MAXLEN_OF_LOGMSG];
                    snprintf(pcLogMsg, _GHM_MAXLEN_OF_LOGMSG,
                             "%s/Host memory allocation failed", __func__);
                    _ghf_logWS_Msg(*pwSet, pcLogMsg); _ghf_wipeWS(pwSet);
                    return(_GHM_HOST_MEMALLOC_ERROR);
                }
                snprintf(pcBldOpts, szBOLen,
                         "%s %s",_GHM_OCLH_COMMON_BUILD_OPTIONS,
                         pcOCLBuildOpts?pcOCLBuildOpts:"");
                pcBldOpts[szBOLen-1]='\0';
#if !defined(__OCLH_COMPILER_ONLY_FLAG) && !defined(__OCLH_LINKER_ONLY_FLAG)
                pwSet->APIErr=clBuildProgram(pwSet->Program, 1u,   &clWrkDev,
                                             pcBldOpts,      NULL, NULL);
                /* TODO: ^[Valg] 326 bytes in 1 blocks are definitely lost in
                 * loss record 1,531 of 1,583 by Intel library
                 * Intel::OpenCL::Framework::ExecutionModule::Initialize */
#endif /* !defined(__OCLH_COMPILER_ONLY_FLAG) &&
          !defined(__OCLH_LINKER_ONLY_FLAG) */
#if defined(__OCLH_COMPILER_ONLY_FLAG) && !defined(__OCLH_LINKER_ONLY_FLAG)
                pwSet->APIErr=
                    clCompileProgram(pwSet->Program, 1ul, &clWrkDev, pcBldOpts,
                                     0u, NULL, NULL, NULL, NULL);
#endif /* defined(__OCLH_COMPILER_ONLY_FLAG) &&
          !defined(__OCLH_LINKER_ONLY_FLAG) */
                cliBuildCLAPIErr=pwSet->APIErr;
                while(_ghf_getWS_BuildStatus(*pwSet)==CL_BUILD_IN_PROGRESS)
                    ;
                {
                    _GHT_LOG TmpLog=_ghf_declLog();
                    /* cl_uint cluNofKernels=0u; */
                    size_t szNofKernels=0ul;
                    cl_build_status clBuildStatus=
                                        _ghf_getWS_BuildStatus(*pwSet);
                    cl_program_binary_type clBinType=0u;
                    _ghf_logWS_BuildInfo(*pwSet, pwSet->Program, clWrkDev,
                                         LogLvl, BuildLogMode);
                    if(BuildLogMode==_GHE_BUILD_LOG_IN_SEPARATED_FILES) {
                        TmpLog=pwSet->Log;
                        int32_t err=0;
                        char pcSepLogFileName[_GHM_MAX_PATH_LENGTH]="\0",
                             pcTextProgramID[_GHM_MAX_PATH_LENGTH]="\0";
                        __ghf_setWS_TextProgramId(*pwSet, pcTextProgramID,
#if defined(__OCLH_BUILDER_FLAG) || \
    defined(__OCLH_COMPILER_ONLY_FLAG) || \
    defined(__OCLH_LINKER_ONLY_FLAG)
                                              pcOutputPrefix);
#else
                                              NULL);
#endif /* defined(__OCLH_BUILDER_FLAG) ||
          defined(__OCLH_COMPILER_ONLY_FLAG) ||
          defined(__OCLH_LINKER_ONLY_FLAG) */
                        snprintf(pcSepLogFileName, _GHM_MAX_PATH_LENGTH,
                                 "%s/%s-%s.log", _GHM_LOG_PATH,
                                 pcTextProgramID,
#if defined(__OCLH_COMPILER_ONLY_FLAG)
                                 "trans"
#endif
#if defined(__OCLH_LINKER_ONLY_FLAG)
                                 "link"
#endif
#if !defined(__OCLH_COMPILER_ONLY_FLAG) && !defined(__OCLH_LINKER_ONLY_FLAG)
                                 "build"
#endif
                                );
                        pwSet->Log=_ghf_declLog();
                        err=_ghf_genrLog(&pwSet->Log, pcSepLogFileName);
                        if(err) { _ghf_wipeLog(&pwSet->Log);
                                  pwSet->Log=TmpLog;
                                  if(pcBldOpts) {
                                      free(pcBldOpts); pcBldOpts=NULL; }
                                  _ghf_wipeWS(pwSet); return(err); }
                        _ghf_logWS_DevInfoShort(*pwSet,
                                                _ghf_getWS_Dev(*pwSet), NULL);
                        _ghf_logWS_BuildInfo(*pwSet, pwSet->Program, clWrkDev,
                                             LogLvl, BuildLogMode);
                    }
#ifdef __OCLH_BUILD_LOG_TO_STDOUT_FLAG
                    {
                        FILE* pTmpFilePtr=pwSet->Log.pfOut;
                        pwSet->Log.pfOut=stdout;
                        _ghf_logWS_DevInfoShort(*pwSet,_ghf_getWS_Dev(*pwSet),
                                                NULL);
                        _ghf_logWS_BuildInfo(*pwSet, pwSet->Program, clWrkDev,
                                             LogLvl, BuildLogMode);
                        pwSet->Log.pfOut=pTmpFilePtr;
                    }
#endif /* __OCLH_BUILD_LOG_TO_STDOUT_FLAG */
                    pwSet->APIErr=cliBuildCLAPIErr;
                    if(pwSet->APIErr) {
                        if(_ghf_isLog_Valid(TmpLog)) {
#if defined(__OCLH_COMPILER_ONLY_FLAG)
                            _ghf_logWS_APIErr(*pwSet, "clCompileProgram");
#endif
#if !defined(__OCLH_COMPILER_ONLY_FLAG) && !defined(__OCLH_LINKER_ONLY_FLAG)
                            _ghf_logWS_APIErr(*pwSet, "clBuildProgram");
#endif
                            _ghf_wipeLog(&pwSet->Log); }
                        pwSet->Log=TmpLog;
                        if(pcBldOpts) { free(pcBldOpts); pcBldOpts=NULL; }
#if defined(__OCLH_COMPILER_ONLY_FLAG)
                        __CLAPI_ERR_ROUTINE_WITH_RET("clCompileProgram")
#endif
#if !defined(__OCLH_COMPILER_ONLY_FLAG) && !defined(__OCLH_LINKER_ONLY_FLAG)
                        __CLAPI_ERR_ROUTINE_WITH_RET("clBuildProgram")
#endif
                    }
                    pwSet->APIErr=
                        clGetProgramBuildInfo(pwSet->Program, clWrkDev,
                                              CL_PROGRAM_BINARY_TYPE,
                                              sizeof(cl_program_binary_type),
                                              &clBinType, NULL);
                    if(pwSet->APIErr) clBinType=CL_PROGRAM_BINARY_TYPE_NONE;
                    if(!clBuildStatus &&
                       (clBinType&CL_PROGRAM_BINARY_TYPE_EXECUTABLE)) {
                        /* another approach to get the number of kernels */
                        /* pwSet->APIErr=
                            clCreateKernelsInProgram(pwSet->Program, 0u, NULL,
                                                     &cluNofKernels); */
                        pwSet->APIErr=
                            clGetProgramInfo(pwSet->Program,
                                             CL_PROGRAM_NUM_KERNELS,
                                             sizeof(size_t), &szNofKernels,
                                             NULL);
                        if(pwSet->APIErr) {
                            if(_ghf_isLog_Valid(TmpLog)) {
                                _ghf_logWS_APIErr(*pwSet, pcLogMsg);
                                _ghf_wipeLog(&pwSet->Log); }
                            pwSet->Log=TmpLog;
                            if(pcBldOpts) { free(pcBldOpts); pcBldOpts=NULL; }
                            __CLAPI_ERR_ROUTINE_WITH_RET(
                                                     "clGetProgramInfo")
                        }
                        pwSet->pKernels=malloc(
                                      sizeof(cl_kernel)*(szNofKernels+1ul));
                        if(!pwSet->pKernels) {
                            char pcLogMsg[_GHM_MAXLEN_OF_LOGMSG];
                            snprintf(pcLogMsg, _GHM_MAXLEN_OF_LOGMSG,
                                     "%s/Host memory allocation failed",
                                     __func__);
                            if(_ghf_isLog_Valid(TmpLog)) {
                                _ghf_logWS_Msg(*pwSet, pcLogMsg);
                                _ghf_wipeLog(&pwSet->Log); }
                            pwSet->Log=TmpLog;
                            _ghf_logWS_Msg(*pwSet, pcLogMsg);
                            if(pcBldOpts) { free(pcBldOpts); pcBldOpts=NULL; }
                            _ghf_wipeWS(pwSet);
                            return(_GHM_HOST_MEMALLOC_ERROR);
                        }
                        pwSet->pKernels[szNofKernels]=NULL;
                        if(szNofKernels) {
                            cl_uint ctrlVal=0u;
                            pwSet->APIErr=
                                clCreateKernelsInProgram(pwSet->Program,
                                                         (cl_uint)szNofKernels,
                                                         pwSet->pKernels,
                                                         &ctrlVal);
                            /* TODO: ^[Valg] 192 (32 direct, 160 indirect) bytes
                             * in 1 blocks are definitely lost
                             * in loss record 1,508 of 1,583 by Intel library
                             * Intel::OpenCL::ClangFE::
                             *     ClangFECompilerGetKernelArgInfoTask::
                             *         GetKernelArgInfo */
                            if(pwSet->APIErr) {
                                if(_ghf_isLog_Valid(TmpLog)) {
                                    _ghf_logWS_APIErr(*pwSet, pcLogMsg);
                                    _ghf_wipeLog(&pwSet->Log); }
                                pwSet->Log=TmpLog;
                                if(pcBldOpts) {
                                    free(pcBldOpts); pcBldOpts=NULL; }
                                __CLAPI_ERR_ROUTINE_WITH_RET(
                                                     "clCreateKernelsInProgram")
                            }
                            if(ctrlVal!=(cl_uint)szNofKernels) {
                                char pcLogMsg[_GHM_MAXLEN_OF_LOGMSG];
                                snprintf(pcLogMsg, _GHM_MAXLEN_OF_LOGMSG,
                                         "%s/The number of created kernels "
                                         " is not equal to the number "
                                         "of all kernels", __func__);
                                if(_ghf_isLog_Valid(TmpLog)) {
                                    _ghf_logWS_Msg(*pwSet, pcLogMsg);
                                    _ghf_wipeLog(&pwSet->Log); }
                                pwSet->Log=TmpLog;
                                _ghf_logWS_Msg(*pwSet, pcLogMsg);
                                if(pcBldOpts) {
                                    free(pcBldOpts); pcBldOpts=NULL; }
                                return(_GHM_CREATING_KERNELS_ERROR);
                            }
                            pwSet->pKernels[szNofKernels]=NULL;
                            /*
                             * TODO: make sort of kernel list for tree search
                             *       for GetKernelByName
                             */
                        }
                    }
                    snprintf(pcLogMsg, _GHM_MAXLEN_OF_LOGMSG,
                             "%lu kernels registered", szNofKernels);
                    _ghf_logWS_Msg(*pwSet, pcLogMsg);
#ifdef __OCLH_BUILD_LOG_TO_STDOUT_FLAG
                    {
                        FILE* pTmpFilePtr=pwSet->Log.pfOut;
                        pwSet->Log.pfOut=stdout;
                        _ghf_logWS_Msg(*pwSet, pcLogMsg);
                        pwSet->Log.pfOut=pTmpFilePtr;
                    }
#endif /* __OCLH_BUILD_LOG_TO_STDOUT_FLAG */
                    if(pwSet->pKernels) {
                        cl_uint i=0u;
                        for(i=0u; pwSet->pKernels[i]; i++) {
                            _ghf_logWS_KerInfo(*pwSet,pwSet->pKernels[i]);
#ifdef __OCLH_BUILD_LOG_TO_STDOUT_FLAG
                            {
                                FILE* pTmpFilePtr=pwSet->Log.pfOut;
                                pwSet->Log.pfOut=stdout;
                                _ghf_logWS_KerInfo(*pwSet,pwSet->pKernels[i]);
                                pwSet->Log.pfOut=pTmpFilePtr;
                            }
#endif /* __OCLH_BUILD_LOG_TO_STDOUT_FLAG */
                        }
                    }
#ifdef __OCLH_BUILD_LOG_TO_STDOUT_FLAG
                    {
                        FILE* pTmpFilePtr=pwSet->Log.pfOut;
                        pwSet->Log.pfOut=stdout;
                        _ghf_logWS_Delim(*pwSet);
                        pwSet->Log.pfOut=pTmpFilePtr;
                    }
#endif /* __OCLH_BUILD_LOG_TO_STDOUT_FLAG */
                    if(BuildLogMode==_GHE_BUILD_LOG_IN_SEPARATED_FILES) {
                        _ghf_wipeLog(&pwSet->Log);
                        pwSet->Log=TmpLog;
                    }
                    if(_ghf_isLog_Valid(TmpLog)) {
                        snprintf(pcLogMsg, _GHM_MAXLEN_OF_LOGMSG,
                             "%lu kernels registered", szNofKernels);
                        _ghf_logWS_Msg(*pwSet, pcLogMsg);
                    }
                }
                if(pcBldOpts) { free(pcBldOpts); pcBldOpts=NULL; }
            }
        } else {
            _ghf_logWS_Msg(*pwSet, "No program sources");
        }
        /* work items and groups*/
        pwSet->szMaxWGsz=_ghf_getWS_MaxWIinWG(*pwSet);
        if(!pwSet->szMaxWGsz) {
            _ghf_wipeWS(pwSet);
            return(_GHM_INVALID_MAXSIZE_OF_WORK_GROUP_ERROR);
        }
        pwSet->szNofCmpUnits=(size_t)_ghf_getWS_MaxCmpUnits(*pwSet);
        if(!pwSet->szNofCmpUnits) {
            _ghf_wipeWS(pwSet); return(_GHM_INVALID_NUMBER_OF_CMP_UNITS_ERROR);
        }
        pwSet->szNofWIinWG=0ul;
        pwSet->szNofAllWI=0ul;
        pwSet->szNofRPThrds=0ul; /* TODO: make detection of
                                          real parallel threads */
        {
            int32_t err=_ghf_flfnWS(*pwSet);
            if(err) { _ghf_wipeWS(pwSet); return(err); }
        }
        pwSet->d_pKerErr=clCreateBuffer(_ghf_getWS_Ctx(*pwSet),
                                        CL_MEM_WRITE_ONLY|CL_MEM_COPY_HOST_PTR,
                                        sizeof(int32_t),
                                        &i32Zero,
                                        &pwSet->APIErr);
        if(pwSet->APIErr) __CLAPI_ERR_ROUTINE_WITH_RET("clCreateBuffer")
        pwSet->KerErr=0;
    }
    pwSet->pszMemBytes=malloc(sizeof(size_t));
    if(!pwSet->pszMemBytes) {
        char pcLogMsg[_GHM_MAXLEN_OF_LOGMSG];
        snprintf(pcLogMsg, _GHM_MAXLEN_OF_LOGMSG,
                 "%s/Host memory allocation failed", __func__);
        _ghf_logWS_Msg(*pwSet,pcLogMsg);
        _ghf_wipeWS(pwSet);
        return(_GHM_HOST_MEMALLOC_ERROR);
    }
    (*pwSet->pszMemBytes)=0ul;
    pwSet->APIErr=0;
    snprintf(pcLogMsg, _GHM_MAXLEN_OF_LOGMSG,
             "Rebuild of workset 0x%lx is finished", (uint64_t)pwSet);
    _ghf_logWS_Hdr(*pwSet, pcLogMsg);
    _ghf_logWS_Delim(*pwSet);
    return(_GHM_OK);
/**/
#undef __CLAPI_CREATE_PROGRAM_ERR_TO_STDOUT
#undef __CLAPI_ERR_ROUTINE_WITH_RET
/**/
}

int32_t _ghf_wipeWS(_GHT_WRKSET* const pwSet) {
/**/
#define __LOG_CLAPI_ERR_ROUTINE(___CLAPI_CALL)                                 \
{ char pcCLAPICall[_GHM_MAXLEN_OF_LOGMSG];                                     \
  snprintf(pcCLAPICall,_GHM_MAXLEN_OF_LOGMSG, "%s/%s", __func__,___CLAPI_CALL);\
  _ghf_logWS_APIErr(*pwSet, pcCLAPICall);                                      \
  i32Res=_GHM_DECONSTRUCTION_OF_WORKSET_ERROR; }
/**/
    int32_t i32Res=_GHM_OK;
    if(pwSet) {
        _GHT_WRKSET* const pRtWS=(_GHT_WRKSET*)pwSet->pwSetAddr;
        cl_context clCtx=NULL;
        char pcLogMsg[_GHM_MAXLEN_OF_LOGMSG];
        if(pwSet->Queue) {
            _ghf_flfnWS(*pwSet);
            if(pRtWS && pRtWS->Queue!=pwSet->Queue) _ghf_flfnWS(*pRtWS);
            clCtx=_ghf_getWS_Ctx(*pwSet);
        }
        if(_ghf_isWS_LogValid(*pwSet)) {
            snprintf(pcLogMsg, _GHM_MAXLEN_OF_LOGMSG,
                     "Deconstruction of workset 0x%lx", (uint64_t)pwSet);
            _ghf_logWS_Msg(*pwSet, pcLogMsg);
        }
        if(pwSet->pfnDatCleaner && pwSet->pvDat) {
            pwSet->pfnDatCleaner(*pwSet);
            pwSet->pvDat=NULL;
            pwSet->pfnDatCleaner=NULL;
        }
        if(pwSet->d_pKerErr) {
            pwSet->APIErr=clReleaseMemObject(pwSet->d_pKerErr);
            if(pwSet->APIErr) __LOG_CLAPI_ERR_ROUTINE("clReleaseMemObject")
            pwSet->d_pKerErr=NULL; if(pRtWS) pRtWS->d_pKerErr=NULL;
        }
        pwSet->KerErr=0; if(pRtWS) pRtWS->KerErr=0; /* ??????? */
        if(pwSet->pKernels) {
            cl_uint i=0u;
            for(i=0u; pwSet->pKernels[i]; i++) {
                pwSet->APIErr=clReleaseKernel(pwSet->pKernels[i]);
                if(pwSet->APIErr) __LOG_CLAPI_ERR_ROUTINE("clReleaseKernel")
            }
            if(pwSet->pKernels) { free(pwSet->pKernels); pwSet->pKernels=NULL; }
            if(pRtWS) pRtWS->pKernels=NULL;
        }
        if(pwSet->Program) {
            pwSet->APIErr=clReleaseProgram(pwSet->Program);
            if(pwSet->APIErr) __LOG_CLAPI_ERR_ROUTINE("clReleaseProgram")
            pwSet->Program=NULL; if(pRtWS) pRtWS->Program=NULL;
        }
        if(_ghf_getWS_Dev(*pwSet)) {
            pwSet->APIErr=clReleaseDevice(_ghf_getWS_Dev(*pwSet));
            if(pwSet->APIErr) __LOG_CLAPI_ERR_ROUTINE("clReleaseDevice")
        }
        if(pwSet->Queue) {
            pwSet->APIErr=clReleaseCommandQueue(pwSet->Queue);
            if(pwSet->APIErr) __LOG_CLAPI_ERR_ROUTINE("clReleaseCommandQueue")
            pwSet->Queue=NULL; if(pRtWS) pRtWS->Queue=NULL;
        }
        if(clCtx) {
            pwSet->APIErr=clReleaseContext(clCtx);
            if(pwSet->APIErr) __LOG_CLAPI_ERR_ROUTINE("clReleaseContext")
        }
        if(pwSet->pszMemBytes) { free(pwSet->pszMemBytes);
                                 pwSet->pszMemBytes=NULL; }
        if(pRtWS) pRtWS->pszMemBytes=NULL;
        pwSet->szNofCmpUnits = 0ul; if(pRtWS) pRtWS->szNofCmpUnits = 0ul;
        pwSet->szNofRPThrds  = 0ul; if(pRtWS) pRtWS->szNofRPThrds  = 0ul;
        pwSet->szMaxWGsz     = 0ul; if(pRtWS) pRtWS->szMaxWGsz     = 0ul;
        pwSet->szNofWIinWG   = 0ul; if(pRtWS) pRtWS->szNofWIinWG   = 0ul;
        pwSet->szNofAllWI    = 0ul; if(pRtWS) pRtWS->szNofAllWI    = 0ul;
        pwSet->APIErr        = 0ul; if(pRtWS) pRtWS->szNofAllWI    = 0ul;
        if(pwSet->i32EigenLogFlag) {
            _ghf_wipeLog(&pwSet->Log); if(pRtWS) pRtWS->Log=_ghf_declLog();
        } else pwSet->Log=_ghf_declLog();
        pwSet->i32EigenLogFlag=0; if(pRtWS) pRtWS->i32EigenLogFlag=0;
        if(pRtWS) {
            pRtWS->pwSetAddr=NULL; pwSet->pwSetAddr=NULL;
            _ghf_wipeWS(pRtWS);
        }
    }
    return(i32Res);
/**/
#undef __LOG_CLAPI_ERR_ROUTINE
/**/
}

int32_t _ghf_isWS_LogValid(const _GHT_WRKSET wSet) {
    return(_ghf_isLog_Valid(wSet.Log));
}

int32_t _ghf_recalcWS_WIWG_1x1(_GHT_WRKSET* const pwSet) {
    pwSet->szNofAllWI=1ul; pwSet->szNofWIinWG=1ul;
    return(_GHM_OK);
}

int32_t _ghf_recalcWS_WIWG_dense(_GHT_WRKSET* const pwSet,
                                 const size_t szNofAllTasks) {
    if((szNofAllTasks/pwSet->szNofCmpUnits)>=pwSet->szMaxWGsz) {
        size_t r=0ul;
        pwSet->szNofWIinWG=pwSet->szMaxWGsz;
        r=szNofAllTasks%(pwSet->szNofWIinWG);
        if(r==0ul) pwSet->szNofAllWI=szNofAllTasks;
        else pwSet->szNofAllWI=szNofAllTasks+(pwSet->szNofWIinWG)-r;
    } else {
        pwSet->szNofWIinWG=szNofAllTasks/pwSet->szNofCmpUnits+1ul;
        pwSet->szNofAllWI=pwSet->szNofWIinWG*pwSet->szNofCmpUnits;
    }
    return(_GHM_OK);
}

int32_t _ghf_flfnWS(_GHT_WRKSET wSet) {
    if(wSet.Queue) {
        wSet.APIErr=clFlush(wSet.Queue);
        if(wSet.APIErr) {
            _ghf_logWS_APIErr(wSet,"clFlush");
            return(_GHM_COMMAND_QUEUE_FLUSH_ERROR);
        }
        wSet.APIErr=clFinish(wSet.Queue);
        if(wSet.APIErr) {
            _ghf_logWS_APIErr(wSet,"clFinish");
            return(_GHM_COMMAND_QUEUE_FINISH_ERROR);
        }
    }
    return(_GHM_OK);
}

cl_device_id _ghf_getWS_Dev(_GHT_WRKSET wSet) {
    cl_device_id clDev=NULL;
    size_t szRes=0ul;
    if(!wSet.Queue) return(_GHM_UNDEFPTR);
    wSet.APIErr=clGetCommandQueueInfo(wSet.Queue, CL_QUEUE_DEVICE,
                                      sizeof(cl_device_id), &clDev, &szRes);
    if(wSet.APIErr) {
        char pcLogMsg[_GHM_MAXLEN_OF_LOGMSG];
        snprintf(pcLogMsg, _GHM_MAXLEN_OF_LOGMSG,
                 "%s/clGetCommandQueueInfo/CL_QUEUE_DEVICE",__func__);
        _ghf_logWS_APIErr(wSet,pcLogMsg);
        clDev=(cl_device_id)_GHM_UNDEFPTR;
    }
    return(clDev);
}

cl_context _ghf_getWS_Ctx(_GHT_WRKSET wSet) {
    cl_context clCtx=NULL;
    size_t szRes=0ul;
    if(!wSet.Queue) return(_GHM_UNDEFPTR);
    wSet.APIErr=clGetCommandQueueInfo(wSet.Queue, CL_QUEUE_CONTEXT,
                                      sizeof(cl_context), &clCtx, &szRes);
    if(wSet.APIErr) {
        char pcLogMsg[_GHM_MAXLEN_OF_LOGMSG];
        snprintf(pcLogMsg, _GHM_MAXLEN_OF_LOGMSG,
                 "%s/clGetCommandQueueInfo/CL_QUEUE_CONTEXT",__func__);
        _ghf_logWS_APIErr(wSet,pcLogMsg);
        clCtx=(cl_context)_GHM_UNDEFPTR;
    }
    return(clCtx);
}

cl_build_status _ghf_getWS_BuildStatus(_GHT_WRKSET wSet) {
    cl_build_status clStatus=0;
    size_t szRes=0ul;
    if(!wSet.Program) return(_GHM_UNDEFBUILDSTATUS);
    wSet.APIErr=clGetProgramBuildInfo(wSet.Program, _ghf_getWS_Dev(wSet),
                                      CL_PROGRAM_BUILD_STATUS,
                                      sizeof(cl_build_status),&clStatus,&szRes);
    if(wSet.APIErr) {
        char pcLogMsg[_GHM_MAXLEN_OF_LOGMSG];
        snprintf(pcLogMsg, _GHM_MAXLEN_OF_LOGMSG,
                 "%s/clGetProgramBuildInfo/CL_PROGRAM_BUILD_STATUS", __func__);
        _ghf_logWS_APIErr(wSet,pcLogMsg);
        clStatus=(cl_build_status) _GHM_UNDEFBUILDSTATUS;
    }
    return(clStatus);
}

cl_kernel _ghf_getWS_KernByName(_GHT_WRKSET wSet,const char* pcKernName) {
    /*
     * TODO: make more effective search
     */
    char* pcCurKernelName=NULL;
    cl_uint i=0u;
    size_t szRes=0ul;
    if(!wSet.pKernels) return(_GHM_UNDEFPTR);
    while(wSet.pKernels[i]) {
        szRes=0ul;
        wSet.APIErr=clGetKernelInfo(wSet.pKernels[i],
                                    CL_KERNEL_FUNCTION_NAME, 0ul, NULL, &szRes);
        if(wSet.APIErr) {
            char pcLogMsg[_GHM_MAXLEN_OF_LOGMSG];
            snprintf(pcLogMsg, _GHM_MAXLEN_OF_LOGMSG,
                     "%s/clGetKernelInfo/CL_KERNEL_FUNCTION_NAME", __func__);
            _ghf_logWS_APIErr(wSet,pcLogMsg);
            return(_GHM_UNDEFPTR);
        }
        pcCurKernelName=malloc(szRes);
        if(!pcCurKernelName) {
            char pcLogMsg[_GHM_MAXLEN_OF_LOGMSG];
            snprintf(pcLogMsg, _GHM_MAXLEN_OF_LOGMSG,
                     "%s/Host memory allocation failed", __func__);
            _ghf_logWS_Msg(wSet,pcLogMsg);
            return(_GHM_UNDEFPTR);
        }
        wSet.APIErr=clGetKernelInfo(wSet.pKernels[i], CL_KERNEL_FUNCTION_NAME,
                                    szRes, pcCurKernelName, &szRes);
        if(wSet.APIErr) {
            char pcLogMsg[_GHM_MAXLEN_OF_LOGMSG];
            snprintf(pcLogMsg, _GHM_MAXLEN_OF_LOGMSG,
                     "%s/clGetKernelInfo/CL_KERNEL_FUNCTION_NAME", __func__);
            _ghf_logWS_APIErr(wSet,pcLogMsg);
            if(pcCurKernelName) { free(pcCurKernelName); pcCurKernelName=NULL; }
            return(_GHM_UNDEFPTR);
        }
        if(!strcmp(pcCurKernelName,pcKernName)) {
            if(pcCurKernelName) { free(pcCurKernelName); pcCurKernelName=NULL; }
            return(wSet.pKernels[i]);
        }
        if(pcCurKernelName) { free(pcCurKernelName); pcCurKernelName=NULL; }
        i++;
    }
    return(_GHM_UNDEFPTR);
}

cl_uint _ghf_getWS_RefCntOfMem(_GHT_WRKSET wSet, const cl_mem clMem) {
    cl_uint refCnt=0u;
    wSet.APIErr=clGetMemObjectInfo(clMem, CL_MEM_REFERENCE_COUNT,
                                   sizeof(cl_uint), &refCnt, NULL);
    if(wSet.APIErr) {
        char pcLogMsg[_GHM_MAXLEN_OF_LOGMSG];
        snprintf(pcLogMsg, _GHM_MAXLEN_OF_LOGMSG,
                 "%s/clGetMemObjectInfo/CL_MEM_REFERENCE_COUNT", __func__);
        _ghf_logWS_APIErr(wSet,pcLogMsg);
        return(0u);
    }
    return(refCnt);
}

size_t _ghf_getWS_MaxWIinWG(_GHT_WRKSET wSet) {
    size_t szInf=0ul;
    if(!wSet.Queue) return(0ul);
    wSet.APIErr=clGetDeviceInfo(_ghf_getWS_Dev(wSet),
                                CL_DEVICE_MAX_WORK_GROUP_SIZE,
                                sizeof(size_t), &szInf, NULL);
    if(wSet.APIErr) {
        char pcLogMsg[_GHM_MAXLEN_OF_LOGMSG];
        snprintf(pcLogMsg, _GHM_MAXLEN_OF_LOGMSG,
                 "%s/clGetDeviceInfo/CL_DEVICE_MAX_WORK_GROUP_SIZE", __func__);
        _ghf_logWS_APIErr(wSet,pcLogMsg);
        return(0ul);
    }
    return(szInf);
}

cl_uint _ghf_getWS_MaxCmpUnits(_GHT_WRKSET wSet) {
    cl_uint cluInf=0u;
    if(!wSet.Queue) return(0u);
    wSet.APIErr=clGetDeviceInfo(_ghf_getWS_Dev(wSet),
                                CL_DEVICE_MAX_COMPUTE_UNITS,
                                sizeof(cl_uint), &cluInf, NULL);
    if(wSet.APIErr) {
        char pcLogMsg[_GHM_MAXLEN_OF_LOGMSG];
        snprintf(pcLogMsg, _GHM_MAXLEN_OF_LOGMSG,
                 "%s/clGetDeviceInfo/CL_DEVICE_MAX_COMPUTE_UNITS", __func__);
        _ghf_logWS_APIErr(wSet,pcLogMsg);
        return(0u);
    }
    return(cluInf);
}

int32_t _ghf_saveWS_ProgramBinaries(_GHT_WRKSET wSet,
                                    char* const pcOutputName,
                                    _GHE_OUTNAME_TYPE OutNameType) {
    char pcLogMsg[_GHM_MAXLEN_OF_LOGMSG]="\0";
    int32_t err=0;
    size_t szRes=0ul,
          *pszOfBins=NULL;
    char** ppcBins=NULL;
    uint64_t u64NofBinsBySizes=0ul,
             u64NofBinsByPtrs=0ul;
    wSet.pvDat=NULL; wSet.pfnDatCleaner=NULL;

    wSet.APIErr=clGetProgramInfo(wSet.Program, CL_PROGRAM_BINARY_SIZES, 0ul,
                                 NULL, &szRes);
    snprintf(pcLogMsg, _GHM_MAXLEN_OF_LOGMSG,
             "%s/clGetProgramInfo/CL_PROGRAM_BINARY_SIZES", __func__);
    if((err=_ghf_wdcChkWS_APIErr (wSet,pcLogMsg, _GHM_NOFL))) return(err);
    pszOfBins=malloc(szRes);
    if(!pszOfBins) return(_GHM_HOST_MEMALLOC_ERROR);

    wSet.APIErr=clGetProgramInfo(wSet.Program, CL_PROGRAM_BINARY_SIZES,
                                 szRes, pszOfBins, NULL);
    snprintf(pcLogMsg, _GHM_MAXLEN_OF_LOGMSG,
             "%s/clGetProgramInfo/CL_PROGRAM_BINARY_SIZES", __func__);
    if((err=_ghf_wdcChkWS_APIErr (wSet,pcLogMsg,_GHM_NOFL))) {
        _ghf_freeHostZ(&pszOfBins); return(err);
    }
    u64NofBinsBySizes=szRes/sizeof(size_t);

    wSet.APIErr=clGetProgramInfo(wSet.Program, CL_PROGRAM_BINARIES, 0ul,
                                 NULL, &szRes);
    snprintf(pcLogMsg, _GHM_MAXLEN_OF_LOGMSG,
             "%s/clGetProgramInfo/CL_PROGRAM_BINARIES", __func__);
    if((err=_ghf_wdcChkWS_APIErr (wSet,pcLogMsg,_GHM_NOFL))) {
        _ghf_freeHostZ(&pszOfBins); return(err);
    }
    ppcBins=malloc(szRes);
    if(!ppcBins) {
        _ghf_freeHostZ(&pszOfBins); return(_GHM_HOST_MEMALLOC_ERROR);
    }
    {
        uint64_t i=0ul;
        int32_t i32AllocErrFlag=0;
        for(i=0ul; i<u64NofBinsBySizes; i++) ppcBins[i]=NULL;
        for(i=0ul; i<u64NofBinsBySizes; i++) ppcBins[i]=malloc(pszOfBins[i]);
        for(i=0ul; i<u64NofBinsBySizes; i++) if(!ppcBins[i]) i32AllocErrFlag=1;
        if(i32AllocErrFlag) {
            for(i=0ul; i<u64NofBinsBySizes; i++) _ghf_freeHostZ(&ppcBins[i]);
             _ghf_freeHostZ(&ppcBins); _ghf_freeHostZ(&pszOfBins);
            return(_GHM_HOST_MEMALLOC_ERROR);
        }
        for(i=0ul; i<u64NofBinsBySizes; i++) ppcBins[i][0]='\0';
    }
    wSet.APIErr=clGetProgramInfo(wSet.Program, CL_PROGRAM_BINARIES, szRes,
                                 ppcBins, NULL);
    snprintf(pcLogMsg, _GHM_MAXLEN_OF_LOGMSG,
             "%s/clGetProgramInfo/CL_PROGRAM_BINARIES", __func__);
    if((err=_ghf_wdcChkWS_APIErr (wSet,pcLogMsg,_GHM_NOFL))) {
        _ghf_freeHostZ(&ppcBins); _ghf_freeHostZ(&pszOfBins);
        return(err);
    }
    u64NofBinsByPtrs=szRes/sizeof(char*);

    if(u64NofBinsByPtrs!=u64NofBinsBySizes) {
        _ghf_freeHostZ(&ppcBins); _ghf_freeHostZ(&pszOfBins);
        return(_GHM_IMPOSSIBLE_SITUATION_ERROR);
    }
    {
        uint64_t i=0ul;
        for(i=0ul; i<u64NofBinsBySizes; i++) {
            FILE* pFile=NULL;
            char pcBinFileName[_GHM_MAX_PATH_LENGTH]="\0",
                 pcTextProgramID[_GHM_MAX_PATH_LENGTH]="\0",
                *pcFileExtension=NULL;
            if(OutNameType==_GHE_VOLATILE_OUTPUT_NAME) {
                __ghf_setWS_TextProgramId(wSet, pcTextProgramID, pcOutputName);
#if defined(__OCLH_COMPILER_ONLY_FLAG)
                pcFileExtension=".clo";
#endif
#if defined(__OCLH_LINKER_ONLY_FLAG)
                {
                    cl_program_binary_type clBinType=0u;
                    size_t szRes=0ul;
                    wSet.APIErr=
                        clGetProgramBuildInfo(wSet.Program,_ghf_getWS_Dev(wSet),
                                              CL_PROGRAM_BINARY_TYPE,
                                              sizeof(cl_program_binary_type),
                                              &clBinType, &szRes);
                    if(wSet.APIErr) clBinType=CL_PROGRAM_BINARY_TYPE_NONE;
                    pcFileExtension=".clout";
                    if(clBinType&CL_PROGRAM_BINARY_TYPE_LIBRARY)
                        pcFileExtension=".clso";
                    if(clBinType&CL_PROGRAM_BINARY_TYPE_EXECUTABLE)
                        pcFileExtension=".clexe";
                }
#endif
#if !defined(__OCLH_COMPILER_ONLY_FLAG) && !defined(__OCLH_LINKER_ONLY_FLAG)
                /* TODO: It is assumed that by default executable binary
                 * is built but actually it is not always so */
                pcFileExtension=".clexe";
#endif
                if(u64NofBinsBySizes>1) {
                    snprintf(pcBinFileName, _GHM_MAX_PATH_LENGTH, "%s/%s%s.%lu",
                            _GHM_LOG_PATH, pcTextProgramID, pcFileExtension, i);
                } else {
                    snprintf(pcBinFileName, _GHM_MAX_PATH_LENGTH, "%s/%s%s",
                            _GHM_LOG_PATH, pcTextProgramID, pcFileExtension);
                }
            }
            if(OutNameType==_GHE_VERBATIM_OUTPUT_NAME) {
                if(pcOutputName && pcOutputName[0]) {
                    if(u64NofBinsBySizes>1) {
                        snprintf(pcBinFileName, _GHM_MAX_PATH_LENGTH,
                                 "%s/%s.%lu", _GHM_LOG_PATH, pcOutputName, i);
                    } else {
                        snprintf(pcBinFileName, _GHM_MAX_PATH_LENGTH, "%s/%s",
                                 _GHM_LOG_PATH, pcOutputName);
                    }
                } else {
                    if(u64NofBinsBySizes>1) {
                        snprintf(pcBinFileName, _GHM_MAX_PATH_LENGTH,
                                 "%s/%s.%lu", _GHM_LOG_PATH, "a.clout", i);
                    } else {
                        snprintf(pcBinFileName, _GHM_MAX_PATH_LENGTH, "%s/%s",
                                 _GHM_LOG_PATH, "a.clout");
                    }
                }
            }
            pFile=fopen(pcBinFileName,"wb");
            if(!pFile) {
                snprintf(pcLogMsg, _GHM_MAXLEN_OF_LOGMSG,
                         "Unable to open the binary file %s", pcBinFileName);
                _ghf_logWS_Msg(wSet,pcLogMsg);
                fprintf(stderr,"%s\n", pcLogMsg);
                perror("");
                { uint64_t i=0ul;
                  for(i=0; i<u64NofBinsBySizes; i++)
                      _ghf_freeHostZ(&ppcBins[i]); }
                _ghf_freeHostZ(&ppcBins); _ghf_freeHostZ(&pszOfBins);
                return(_GHM_OPEN_FILE_ERROR);
            }
            if(fwrite(ppcBins[i],pszOfBins[i],1,pFile)!=1) {
                snprintf(pcLogMsg, _GHM_MAXLEN_OF_LOGMSG,
                         "Unable to write the binary file %s", pcBinFileName);
                _ghf_logWS_Msg(wSet,pcLogMsg);
                fprintf(stderr,"%s\n", pcLogMsg);
                perror("");
                { uint64_t i=0ul;
                  for(i=0; i<u64NofBinsBySizes; i++)
                      _ghf_freeHostZ(&ppcBins[i]); }
                _ghf_freeHostZ(&ppcBins); _ghf_freeHostZ(&pszOfBins);
                fclose(pFile);
                return(_GHM_WRITE_FILE_ERROR);
            }
            fclose(pFile);
            snprintf(pcLogMsg, _GHM_MAXLEN_OF_LOGMSG,
                     "The program binary have been written to %s",
                     pcBinFileName);
            _ghf_logWS_Msg(wSet,pcLogMsg);
        }
    }
    { uint64_t i=0ul;
      for(i=0ul; i<u64NofBinsBySizes; i++) _ghf_freeHostZ(&ppcBins[i]); }
    _ghf_freeHostZ(&ppcBins);
    _ghf_freeHostZ(&pszOfBins);
    return(_GHM_OK);
}

int32_t _ghf_wdcChkWS_APIErr(_GHT_WRKSET wSet,
                             const char* const pcAPICall,
                             const int32_t i32FlashFlag) {
    if(wSet.APIErr) {
        _ghf_logWS_APIErr(wSet,pcAPICall);
        if(wSet.pvDat) wSet.pfnDatCleaner(wSet);
        _ghf_flfnWS(wSet);
        return(wSet.APIErr);
    } else {
        if(i32FlashFlag) {
            if(_ghf_flfnWS(wSet)) {
                if(wSet.pvDat) wSet.pfnDatCleaner(wSet);
                return(wSet.APIErr);
            }
        }
    }
    return(_GHM_OK);
}

int32_t _ghf_wdcChkWS_KerErr(_GHT_WRKSET wSet, const cl_kernel clKer) {
    int32_t err=_GHM_OK;
    char pcLogMsg[_GHM_MAXLEN_OF_LOGMSG];
    snprintf(pcLogMsg, _GHM_MAXLEN_OF_LOGMSG,
             "%s/clEnqueueReadBuffer", __func__);
    wSet.APIErr=clEnqueueReadBuffer(wSet.Queue, wSet.d_pKerErr, CL_TRUE, 0ul,
                                    sizeof(int32_t),&wSet.KerErr,0u,NULL,NULL);
    if((err=_ghf_wdcChkWS_APIErr(wSet,pcLogMsg,_GHM_FL))) return(err);
    if(wSet.KerErr) {
        char pcKerName[_GHM_MAXLEN_OF_LOGMSG];
        { /* get kernel name */
            size_t szRes=0ul;
            snprintf(pcLogMsg, _GHM_MAXLEN_OF_LOGMSG,
                     "%s/clGetKernelInfo", __func__);
            wSet.APIErr=clGetKernelInfo(clKer, CL_KERNEL_FUNCTION_NAME, 0ul,
                                        NULL, &szRes);
            if((err=_ghf_wdcChkWS_APIErr(wSet,pcLogMsg,_GHM_FL))) return(err);
            wSet.APIErr=clGetKernelInfo(clKer, CL_KERNEL_FUNCTION_NAME,
                                        (szRes>_GHM_MAXLEN_OF_LOGMSG)?
                                            _GHM_MAXLEN_OF_LOGMSG : szRes,
                                        pcKerName, &szRes);
            if((err=_ghf_wdcChkWS_APIErr(wSet,pcLogMsg,_GHM_FL))) return(err);
        }
        snprintf(pcLogMsg, _GHM_MAXLEN_OF_LOGMSG,
                 "Error %d in kernel %s", wSet.KerErr, pcKerName);
        _ghf_logWS_Msg(wSet,pcLogMsg);
        if(wSet.pvDat) wSet.pfnDatCleaner(wSet);
        _ghf_flfnWS(wSet);
        return(wSet.KerErr);
    }
    return(wSet.KerErr);
}

int32_t _ghf_wdcSetWS_KerErrToZero(_GHT_WRKSET wSet) {
    const int32_t i32Zero=0;
    int32_t err=_GHM_OK;
    char pcLogMsg[_GHM_MAXLEN_OF_LOGMSG];
    snprintf(pcLogMsg, _GHM_MAXLEN_OF_LOGMSG,
             "%s/clEnqueueWriteBuffer", __func__);
    wSet.APIErr=clEnqueueWriteBuffer(wSet.Queue, wSet.d_pKerErr, CL_TRUE, 0ul,
                                     sizeof(int32_t), &i32Zero, 0u, NULL, NULL);
    if((err=_ghf_wdcChkWS_APIErr(wSet,pcLogMsg,_GHM_FL))) return(err);
    return(_GHM_OK);
}

/**/
#define __GDM_SET_KERNEL_ARGS_BY_VA_LIST_FROM_NAME_W_RET(___GDM_SKA_MAC_Arg) \
{ \
    cl_uint cluNofKerArgs=0u, i=0u; \
    size_t szRes=0ul; \
    char pcLogMsg[_GHM_MAXLEN_OF_LOGMSG]; \
    va_list volArgList; \
    wSet.APIErr=clGetKernelInfo(kr, CL_KERNEL_NUM_ARGS, \
                                sizeof(cl_uint), &cluNofKerArgs, &szRes); \
    snprintf(pcLogMsg, _GHM_MAXLEN_OF_LOGMSG, \
             "%s/clGetKernelInfo/CL_KERNEL_NUM_ARGS", __func__); \
    if(_ghf_wdcChkWS_APIErr(wSet,pcLogMsg,_GHM_NOFL)) return(wSet.APIErr); \
    va_start(volArgList,___GDM_SKA_MAC_Arg); \
    snprintf(pcLogMsg, _GHM_MAXLEN_OF_LOGMSG,"%s/clSetKernelArg",__func__); \
    for(i=0;i<cluNofKerArgs;i++) { \
        size_t sz=va_arg(volArgList,size_t); \
        const void* addr=va_arg(volArgList,const void*); \
        wSet.APIErr=clSetKernelArg(kr,i,sz,addr); \
        if(_ghf_wdcChkWS_APIErr(wSet,pcLogMsg,_GHM_NOFL)) return(wSet.APIErr); \
    } \
    va_end(volArgList); \
}
/**/

int32_t _ghf_wdcSetWS_KernArgs(_GHT_WRKSET wSet,const cl_kernel kr,...) {
    __GDM_SET_KERNEL_ARGS_BY_VA_LIST_FROM_NAME_W_RET(kr);
    return(_GHM_OK);
}
int32_t _ghf_wdcBExecWS_Kern(_GHT_WRKSET wSet,
                            const char* const pcKernName, ...) {
    int32_t err=0;
    const cl_kernel kr=_ghf_getWS_KernByName(wSet,pcKernName);
    if(kr==_GHM_UNDEFPTR) {
        if(wSet.pvDat) wSet.pfnDatCleaner(wSet);
        return(_GHM_UNAVALABLE_KERNEL_ERROR);
    }
    if((err=_ghf_wdcSetWS_KerErrToZero(wSet))) return(err);
    __GDM_SET_KERNEL_ARGS_BY_VA_LIST_FROM_NAME_W_RET(pcKernName);
    wSet.APIErr=clEnqueueNDRangeKernel(wSet.Queue,kr,1u,0ul,
                                       &wSet.szNofAllWI,&wSet.szNofWIinWG,
                                       0u,NULL,NULL);
    if((err=_ghf_wdcChkWS_APIErr(wSet,"clEnqueueNDRangeKernel",_GHM_FL)))
                                                                    return(err);
    if((err=_ghf_wdcChkWS_KerErr(wSet,kr))) return(err);
    return(_GHM_OK);
}
/**/
#undef __GDM_SET_KERNEL_ARGS_BY_VA_LIST_FROM_NAME_W_RET
/**/

_GHT_DEVLIST_DESC _ghf_declDevLstDesc(void) {
    _GHT_DEVLIST_DESC dld={ .pcluIdxs=NULL, .pcWC=NULL };
    return(dld);
}

int32_t _ghf_isDevLstDesc_Valid(const _GHT_DEVLIST_DESC DevLstDesc) {
    if(DevLstDesc.pcluIdxs &&  DevLstDesc.pcWC)        return(0);
    if(DevLstDesc.pcluIdxs &&  DevLstDesc.pcluIdxs[0]==(cl_uint)(-1)) return(0);
    if(DevLstDesc.pcWC     && !DevLstDesc.pcWC)        return(0);
    return(1);
}

int32_t _ghf_addDevIdxToDevLstDesc(   _GHT_DEVLIST_DESC* const pDevLstDesc,
                                      const cl_uint            cluIdx) {
    if(!pDevLstDesc->pcluIdxs) {
        pDevLstDesc->pcluIdxs=malloc(sizeof(cl_uint));
        if(!pDevLstDesc->pcluIdxs) return(_GHM_HOST_MEMALLOC_ERROR);
        pDevLstDesc->pcluIdxs[0]=(cl_uint)(-1);
    }
    _ghf_freeHostZ(&pDevLstDesc->pcWC);
    {
        uint64_t i=0ul;
        cl_uint* pTmp=NULL;
        /* TODO: make faster */
        for(i=0ul; pDevLstDesc->pcluIdxs[i]!=(cl_uint)(-1); i++ )
            if(pDevLstDesc->pcluIdxs[i]==cluIdx) return(_GHM_OK);
        pTmp=realloc(pDevLstDesc->pcluIdxs,i+1);
        if(!pTmp) return(_GHM_HOST_MEMALLOC_ERROR);
        pDevLstDesc->pcluIdxs=pTmp;
        pDevLstDesc->pcluIdxs[i]=cluIdx;
        pDevLstDesc->pcluIdxs[i+1]=(cl_uint)(-1);
    }
    return(_GHM_OK);
}

int32_t _ghf_setDevLstDesc_DevNameWC(_GHT_DEVLIST_DESC* const pDevLstDesc,
                                     const char* const pcWC) {
    char* pTmp=NULL;
    uint64_t i=0;
    while(pcWC[i]) i++;
    pTmp=realloc(pDevLstDesc->pcWC,i+1ul);
    if(!pTmp) return(_GHM_HOST_MEMALLOC_ERROR);
    pDevLstDesc->pcWC=pTmp;
    strncpy(pDevLstDesc->pcWC,pcWC,i);
    pDevLstDesc->pcWC[i]='\0';
    _ghf_freeHostZ(&pDevLstDesc->pcluIdxs);
    return(_GHM_OK);
}

int32_t _ghf_wipeDevLstDesc(          _GHT_DEVLIST_DESC* const pDevLstDesc) {
    _ghf_freeHostZ(&pDevLstDesc->pcluIdxs);
    _ghf_freeHostZ(&pDevLstDesc->pcWC);
    return(0);
}

_GHT_AWSS _ghf_declAWSs(const _GHE_LOGLVL  LogLvl,
                        const _GHE_BUILD_LOG_MODE BuildLogMode) {
    _GHT_AWSS AWSs={ .pWSet=NULL, .u64NofWSs=0ul, .Log=_ghf_declLog(),
                     .pvCLProgramSource=NULL,
                     .SourceType=_GHE_HL_LISTINGS,
                     .LogLevel=LogLvl, .BuildLogMode=BuildLogMode };
    return(AWSs);
}

int32_t _ghf_genrAWSs(      _GHT_AWSS*   const pAWSs,
                      const _GHT_LOG           Log,
                      const _GHT_DEVLIST_DESC  DevLstDesc,
                            char**       const ppcCLProgramSources,
                      const _GHE_SRCTYPE       SourceType,
                      const char*        const OCLBuildOpts
#if defined(__OCLH_BUILDER_FLAG) || \
    defined(__OCLH_COMPILER_ONLY_FLAG) || \
    defined(__OCLH_LINKER_ONLY_FLAG)
                    ,       char*        const pcOutputPrefix
#endif /* defined(__OCLH_BUILDER_FLAG) ||
          defined(__OCLH_COMPILER_ONLY_FLAG) ||
          defined(__OCLH_LINKER_ONLY_FLAG) */
                                                      ) {
    char pcLogMsg[_GHM_MAXLEN_OF_LOGMSG];
    _GHT_WRKSET tmpWS=_ghf_declWS();
    cl_device_id* DevLst=NULL;
    int32_t err=0;
    _GHT_DEVLIST_DESC TmpDevLstDesc=_ghf_declDevLstDesc();
    if(!_ghf_isLog_Valid(Log)) return(_GHM_INCONSISTED_LOG_ERROR);
    if(!_ghf_isDevLstDesc_Valid(DevLstDesc))
                               return(_GHM_INCONSISTED_DEV_DESCRIPTION_ERROR);
    pAWSs->Log=Log;
    tmpWS.Log=Log; tmpWS.pwSetAddr=&tmpWS;
    _ghf_wipeAWSs(pAWSs);
    if(pAWSs->LogLevel>=_GHE_LOG_ALL) {
        snprintf(pcLogMsg,_GHM_MAXLEN_OF_LOGMSG,
                 "Temp workset 0x%lx created",(uint64_t)&tmpWS );
        _ghf_logWS_Msg(tmpWS,pcLogMsg);
    }
    if((err=_ghf_buildDevList(tmpWS,&DevLst))) {
        if(err==_GHM_NO_PLATFORMS_AVAILABLE)
            _ghf_logWS_Msg(tmpWS,
                           "There is no an accessible OpenCL platforms here");
        if(err==_GHM_NO_DEVICES_AVAILABLE)
            _ghf_logWS_Msg(tmpWS,
                           "There is no an accessible OpenCL device here");
        _ghf_wipeAWSs(pAWSs);
        _ghf_wipeDevList(&DevLst);
        _ghf_wipeWS(&tmpWS);
        return(err);
    }
    {
        cl_uint i=0u;
        while(DevLst[i]) i++;
        snprintf(pcLogMsg, _GHM_MAXLEN_OF_LOGMSG, "Devices available: %u", i);
        _ghf_logWS_Msg(tmpWS,pcLogMsg);
        _ghf_logWS_Delim(tmpWS);
        for(i=0u; DevLst[i]; i++)
            if(pAWSs->LogLevel>=_GHE_LOG_KEY_PARS) {
                snprintf(pcLogMsg,_GHM_MAXLEN_OF_LOGMSG,
                         "+++++++ Device 0x%lx description +++++++",
                         (uint64_t)DevLst[i]);
                _ghf_logWS_Msg(tmpWS, pcLogMsg);
                snprintf(pcLogMsg, _GHM_MAXLEN_OF_LOGMSG,
                         "dev_0x%04lx | Device index: %u",
                         __GHM_U64STRIPTO2B(DevLst[i]),i);
                _ghf_logWS_Msg(tmpWS, pcLogMsg);
                _ghf_logWS_DevInfo(tmpWS,DevLst[i],pAWSs->LogLevel);
                snprintf(pcLogMsg,_GHM_MAXLEN_OF_LOGMSG,
                         "------- End of device 0x%lx description -------",
                         (uint64_t)DevLst[i]);
                _ghf_logWS_Msg(tmpWS, pcLogMsg);
            }
    }
    {
        cl_uint i=0u,
                cluNofDevs=0u,
                cluNofWS=0u;
        err=0;
        while(DevLst[cluNofDevs]) cluNofDevs++;
        if(!DevLstDesc.pcluIdxs && !DevLstDesc.pcWC) { cluNofWS=cluNofDevs; }
        else {
            if(DevLstDesc.pcluIdxs) {
                cl_uint j=0u;
                while(DevLstDesc.pcluIdxs[j]!=(cl_uint)(-1)) {
                    if(DevLstDesc.pcluIdxs[j]<cluNofDevs) { cluNofWS++; }
                    j++;
                }
            } else { /* подразумевается валидное описание устройств */
                cl_uint j=0u;
                while(DevLst[j]) {
                    char* pcDevName=_ghf_getDevInf_charptr(DevLst[j],
                                                           CL_DEVICE_NAME,
                                                           &tmpWS.APIErr);
                    if(__ghf_cmpCharPtrAndWC(pcDevName,DevLstDesc.pcWC,
                                             __GHE_CASE_SENSITIVE)) {
                        err=_ghf_addDevIdxToDevLstDesc(&TmpDevLstDesc,j);
                        if(err) {
                            if(pcDevName) { free(pcDevName);  pcDevName =NULL; }
                            _ghf_wipeDevLstDesc(&TmpDevLstDesc);
                            _ghf_wipeAWSs(pAWSs);
                            _ghf_wipeDevList(&DevLst);
                            _ghf_wipeWS(&tmpWS);
                            return(err);
                        }
                        cluNofWS++;
                    }
                    if(pcDevName) { free(pcDevName);  pcDevName=NULL; }
                    j++;
                }
            }
        }
        err=0;
        pAWSs->u64NofWSs=0;
        pAWSs->pWSet=malloc(sizeof(_GHT_WRKSET)*cluNofWS);
        if(!pAWSs->pWSet) {
            _ghf_wipeDevLstDesc(&TmpDevLstDesc);
            _ghf_wipeAWSs(pAWSs);
            _ghf_wipeDevList(&DevLst);
            _ghf_wipeWS(&tmpWS);
            return(_GHM_HOST_MEMALLOC_ERROR);
        }
        for(i=0u; i<cluNofWS; i++) {
            pAWSs->pWSet[pAWSs->u64NofWSs]=_ghf_declWS();
            err=_ghf_genrWS(&pAWSs->pWSet[pAWSs->u64NofWSs], Log, 0,
                            (TmpDevLstDesc.pcluIdxs)?
                                DevLst[TmpDevLstDesc.pcluIdxs[i]]:
                                DevLst[DevLstDesc.pcluIdxs?
                                           DevLstDesc.pcluIdxs[i]:i],
                            ppcCLProgramSources, SourceType, OCLBuildOpts,
                            pAWSs->LogLevel, pAWSs->BuildLogMode
#if defined(__OCLH_BUILDER_FLAG) || \
    defined(__OCLH_COMPILER_ONLY_FLAG) || \
    defined(__OCLH_LINKER_ONLY_FLAG)
                          , pcOutputPrefix
#endif /* defined(__OCLH_BUILDER_FLAG) ||
          defined(__OCLH_COMPILER_ONLY_FLAG) ||
          defined(__OCLH_LINKER_ONLY_FLAG) */
                                          );
            if(err) _ghf_wipeWS(&pAWSs->pWSet[pAWSs->u64NofWSs]);
            else pAWSs->u64NofWSs++;
        }
        if(!pAWSs->u64NofWSs) {
            _ghf_wipeDevLstDesc(&TmpDevLstDesc);
            _ghf_wipeAWSs(pAWSs);
            _ghf_wipeDevList(&DevLst);
            _ghf_wipeWS(&tmpWS);
            return(err);
        }
    }
    _ghf_wipeDevLstDesc(&TmpDevLstDesc);
    _ghf_wipeDevList(&DevLst);
    if(pAWSs->LogLevel>=_GHE_LOG_ALL) {
        snprintf(pcLogMsg, _GHM_MAXLEN_OF_LOGMSG,
                 "Temp workset 0x%lx destroyed", (uint64_t)&tmpWS);
        _ghf_logWS_Msg(tmpWS, pcLogMsg);
    }
    _ghf_wipeWS(&tmpWS);
    return(_GHM_OK);
}

int32_t _ghf_wipeAWSs(_GHT_AWSS* const pAWSs) {
    if(pAWSs) {
        uint64_t i=0ul;
        for(i=0ul; i<pAWSs->u64NofWSs; i++)
            _ghf_wipeWS(&pAWSs->pWSet[i]);
        if(pAWSs->pWSet) { free(pAWSs->pWSet); pAWSs->pWSet=NULL; }
        pAWSs->u64NofWSs=0ul;
    }
    return(_GHM_OK);
}

int32_t _ghf_buildDevList(_GHT_WRKSET wSet, cl_device_id** ppDevLst) {
    _ghf_wipeDevList(ppDevLst);
    cl_uint cluNofPlatfs=0u;
    cl_platform_id* pPlatform=NULL;
    cl_uint i=0u;
    /* get platforms */
    wSet.APIErr=clGetPlatformIDs(0u, NULL, &cluNofPlatfs);
    /* 278 (240 direct, 38 indirect) bytes in 1 blocks are definitely lost
     * in loss record 1,525 of 1,583 by Intel library
     * Intel::OpenCL::Framework::ExecutionModule::Initialize
     * Intel::OpenCL::Framework::Device::CreateAndInitAllDevicesOfDeviceType */
    if(wSet.APIErr) {
        char pcLogMsg[_GHM_MAXLEN_OF_LOGMSG];
        snprintf(pcLogMsg, _GHM_MAXLEN_OF_LOGMSG,
                 "%s/clGetPlatformIDs", __func__);
        _ghf_logWS_APIErr(wSet,pcLogMsg);
        _ghf_wipeDevList(ppDevLst);
        return(wSet.APIErr);
    }
    if(!cluNofPlatfs) { _ghf_wipeDevList(ppDevLst);
                        return(_GHM_NO_PLATFORMS_AVAILABLE); }
    pPlatform=malloc(sizeof(cl_platform_id)*cluNofPlatfs);
    if(!pPlatform) {
        char pcLogMsg[_GHM_MAXLEN_OF_LOGMSG];
        snprintf(pcLogMsg, _GHM_MAXLEN_OF_LOGMSG,
                 "%s/Host memory allocation failed", __func__);
        _ghf_logWS_Msg(wSet,pcLogMsg);
        _ghf_wipeDevList(ppDevLst);
        return(_GHM_HOST_MEMALLOC_ERROR);
    }
    wSet.APIErr=clGetPlatformIDs(cluNofPlatfs, pPlatform, NULL);
    if(wSet.APIErr) {
        char pcLogMsg[_GHM_MAXLEN_OF_LOGMSG];
        snprintf(pcLogMsg, _GHM_MAXLEN_OF_LOGMSG,
                 "%s/clGetPlatformIDs", __func__);
        _ghf_logWS_APIErr(wSet,pcLogMsg);
        if(pPlatform) { free(pPlatform); pPlatform=NULL; }
        _ghf_wipeDevList(ppDevLst);
        return(wSet.APIErr);
    }
    /* platforms have been got */
    { /* devices */
        uintptr_t ptrOff=0ul;
        cl_uint cluCurPlatfDevsCount = 0u,
                cluCurPlatfNum       = 0u,
                cluNofDevs           = 0u,
               *pcluNumByPlatf       = NULL;
        pcluNumByPlatf=malloc(sizeof(cl_uint)*cluNofPlatfs);
        if(!pcluNumByPlatf) {
            char pcLogMsg[_GHM_MAXLEN_OF_LOGMSG];
            snprintf(pcLogMsg, _GHM_MAXLEN_OF_LOGMSG,
                     "%s/Host memory allocation failed", __func__);
            _ghf_logWS_Msg(wSet,pcLogMsg);
            if(pPlatform) { free(pPlatform); pPlatform=NULL; }
            _ghf_wipeDevList(ppDevLst);
            return(_GHM_HOST_MEMALLOC_ERROR);
        }
        for(i=0u; i<cluNofPlatfs; i++) {
            wSet.APIErr=clGetDeviceIDs(pPlatform[i], CL_DEVICE_TYPE_ALL, 0u,
                                       NULL, &cluCurPlatfDevsCount);
            if(wSet.APIErr) {
                char pcLogMsg[_GHM_MAXLEN_OF_LOGMSG];
                snprintf(pcLogMsg, _GHM_MAXLEN_OF_LOGMSG,
                         "%s/clGetDeviceIDs", __func__);
                _ghf_logWS_APIErr(wSet,pcLogMsg);
                if(pcluNumByPlatf) { free(pcluNumByPlatf);pcluNumByPlatf=NULL; }
                if(pPlatform) { free(pPlatform); pPlatform=NULL; }
                _ghf_wipeDevList(ppDevLst);
                return(wSet.APIErr);
            }
            pcluNumByPlatf[i]=cluCurPlatfDevsCount;
            cluNofDevs+=cluCurPlatfDevsCount;
        }
        if(!cluNofDevs) {
            if(pcluNumByPlatf) { free(pcluNumByPlatf); pcluNumByPlatf=NULL; }
            if(pPlatform) { free(pPlatform); pPlatform=NULL; }
            _ghf_wipeDevList(ppDevLst);
            return(_GHM_NO_DEVICES_AVAILABLE);
        }
        *ppDevLst=malloc(sizeof(cl_device_id)*(cluNofDevs+1ul));
        if(!*ppDevLst) {
            char pcLogMsg[_GHM_MAXLEN_OF_LOGMSG];
            snprintf(pcLogMsg, _GHM_MAXLEN_OF_LOGMSG,
                     "%s/Host memory allocation failed", __func__);
            _ghf_logWS_Msg(wSet,pcLogMsg);
            if(pcluNumByPlatf) { free(pcluNumByPlatf); pcluNumByPlatf=NULL; }
            if(pPlatform) { free(pPlatform); pPlatform=NULL; }
            _ghf_wipeDevList(ppDevLst);
            return(_GHM_HOST_MEMALLOC_ERROR);
        }
        for(i=0u; i<cluNofPlatfs; i++){
            wSet.APIErr=clGetDeviceIDs(pPlatform[i], CL_DEVICE_TYPE_ALL,
                                       pcluNumByPlatf[i],
                                       &((*ppDevLst)[ptrOff]),&cluCurPlatfNum);
            (*ppDevLst)[ptrOff+cluCurPlatfNum]=NULL;
            if(wSet.APIErr) {
                char pcLogMsg[_GHM_MAXLEN_OF_LOGMSG];
                snprintf(pcLogMsg, _GHM_MAXLEN_OF_LOGMSG,
                         "%s/clGetDeviceIDs", __func__);
                _ghf_logWS_APIErr(wSet,pcLogMsg);
                if(pcluNumByPlatf) { free(pcluNumByPlatf);pcluNumByPlatf=NULL; }
                if(pPlatform) { free(pPlatform); pPlatform=NULL; }
                _ghf_wipeDevList(ppDevLst); return(wSet.APIErr);
            }
            ptrOff+=cluCurPlatfNum;
        }
        if(pcluNumByPlatf) { free(pcluNumByPlatf); pcluNumByPlatf=NULL; }
    }
    if(pPlatform) { free(pPlatform); pPlatform=NULL; }
    return(_GHM_OK);
}

int32_t _ghf_wipeDevList(cl_device_id** const ppDevLst) {
    if(ppDevLst && *ppDevLst) { free(*ppDevLst); *ppDevLst=NULL; }
    return(_GHM_OK);
}



int32_t __ghf_setWS_TextProgramId(_GHT_WRKSET wSet,
                                  char* const pcDst,
                                  char* const pcOutputPrefix) {
    char pcDefProgName[64]="\0",
         pcDefDevName[64]="\0",
        *pcApiDevName=_ghf_getDevInf_charptr(_ghf_getWS_Dev(wSet),
                                             CL_DEVICE_NAME, &wSet.APIErr);
    snprintf(pcDefProgName, 64ul,
             "program_0x%04lx", __GHM_U64STRIPTO2B((uint64_t)wSet.Program));
    snprintf(pcDefDevName, 64ul,
             "dev_0x%04lx", __GHM_U64STRIPTO2B((uint64_t)_ghf_getWS_Dev(wSet)));
    if(pcOutputPrefix) {
        __ghf_removePreNPostSpacesFromCharPtr(pcOutputPrefix);
        __ghf_replaceSpacesWithUnderscoreInCharPtr(pcOutputPrefix);
    }
    if(pcApiDevName) {
        __ghf_removePreNPostSpacesFromCharPtr(pcApiDevName);
        __ghf_replaceIllegalFSCharsInCharPtr(pcApiDevName);
        __ghf_replaceSpacesWithUnderscoreInCharPtr(pcApiDevName);
    }
    snprintf(pcDst,_GHM_MAX_PATH_LENGTH,"%s/%s-%s",_GHM_LOG_PATH,
             (pcOutputPrefix)?pcOutputPrefix:pcDefProgName,
             (pcApiDevName)?pcApiDevName:pcDefDevName);
    if(pcApiDevName) { free(pcApiDevName); pcApiDevName=NULL; }
    return(_GHM_OK);
}
