_ghf_wdcSetWS_KernArgs, _ghf_wdcExecWS_Kern added

src/include_h/oclh_h_ws_base.h

@@ -6,6 +6,7 @@
 #define OCLH_H_WS_BASE_H_ 1
 #include <sys/types.h>
 #include <stdlib.h>
+#include <stdarg.h>
 #include <CL/opencl.h>
 #include <oclh_h_base_log.h>
 
@@ -77,8 +78,9 @@ int32_t _ghf_genrWS(      _GHT_WRKSET*  const pwSet,
 int32_t _ghf_wipeWS(_GHT_WRKSET* const pwSet);
 
 int32_t _ghf_isWS_LogValid(const _GHT_WRKSET wSet);
-int32_t _ghf_recalcWS_WIWG(_GHT_WRKSET* const pwSet,
-                           const size_t szNofAllTasks);
+int32_t _ghf_recalcWS_WIWG_1x1(  _GHT_WRKSET* const pwSet);
+int32_t _ghf_recalcWS_WIWG_dense(_GHT_WRKSET* const pwSet,
+                           const size_t             szNofAllTasks);
 int32_t _ghf_flfnWS(_GHT_WRKSET wSet);
 
 cl_device_id    _ghf_getWS_Dev(        _GHT_WRKSET wSet);
@@ -97,6 +99,10 @@ int32_t         _ghf_wdcChkWS_APIErr(      _GHT_WRKSET wSet,
 int32_t         _ghf_wdcChkWS_KerErr(      _GHT_WRKSET wSet,
                                      const cl_kernel   clKer);
 int32_t         _ghf_wdcSetWS_KerErrToZero(_GHT_WRKSET wSet);
+int32_t         _ghf_wdcSetWS_KernArgs(    _GHT_WRKSET wSet,
+                                     const cl_kernel kr, ...);
+int32_t         _ghf_wdcExecWS_Kern(       _GHT_WRKSET wSet,
+                                     const char* const pcKernName, ...);
 
 #pragma pack(push,1)
 typedef struct _GHT_LIST_OF_DEVICES_DESCRIPTION {

src/oclh_h_ws_base.c

@@ -664,8 +664,13 @@ int32_t _ghf_isWS_LogValid(const _GHT_WRKSET wSet) {
     return(_ghf_isLog_Valid(wSet.Log));
 }
 
-int32_t _ghf_recalcWS_WIWG(_GHT_WRKSET* const pwSet,
-                           const size_t szNofAllTasks) {
+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;
@@ -1069,6 +1074,56 @@ int32_t _ghf_wdcSetWS_KerErrToZero(_GHT_WRKSET wSet) {
     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_wdcExecWS_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);