builder fairly usable

Makefile

@@ -0,0 +1,80 @@
+-include makefile.init
+#-include makefile.defs
+RM := rm -rf
+CC := gcc
+LINKER := gcc
+CFLAGS := -std=c11 -pedantic -Wall -Wextra -Wconversion
+LDFLAGS := -L/usr/local/cuda/lib64 -L/opt/intel/opencl/lib64
+OUTPUTFLAGS := -c -fmessage-length=0
+MISCFLAGS := -pthread
+INCLUDES := -I"src/include_h/" -I"src/include_hd/" \
+            -I"/usr/local/cuda/include"
+DEFS :=
+# for debug
+MODEFLAGS := -O0 -g3
+TARGETDIR := build/Debug
+# for release
+### MODEFLAGS := -O3
+### TARGETDIR := build/Release
+# All of the sources participating in the build are defined here
+# -include sources.mk
+OBJS :=
+C_DEPS :=
+# Every subdirectory with source files must be described here
+SRCDIR := src
+# -include subdir.mk
+# Add inputs and outputs from these tool invocations to the build variables
+C_SRCS += $(wildcard $(SRCDIR)/*.c)
+#OBJS += $(wildcard $(TARGETDIR)/*.o)
+OBJS += $(patsubst $(SRCDIR)/%.c,$(TARGETDIR)/%.o,$(C_SRCS))
+C_DEPS += $(wildcard $(TARGETDIR)/*.d)
+
+OCLH_BUILDER_NAME  := oclh_br
+OCLH_COMPILER_NAME := oclh_cr
+OCLH_LINKER_NAME   := oclh_lr
+OCLH_LIBRARY_NAME  := liboclh.so
+
+# Each subdirectory must supply rules for building sources it contributes
+$(TARGETDIR)/%.o: $(SRCDIR)/%.c
+	@echo 'Building file: $<'
+	@echo 'Invoking: GCC C Compiler'
+	mkdir -p $(TARGETDIR); \
+	$(CC) $(INCLUDES) $(MODEFLAGS) $(CFLAGS) $(DEFS) $(OUTPUTFLAGS) \
+	-MMD -MP -MF"$(@:%.o=%.d)" -MT"$(@)" \
+	-o "$@" "$<"
+	@echo 'Finished building: $<'
+	@echo ' '
+# -include objects.mk
+USER_OBJS :=
+LIBS := -lOpenCL
+
+ifneq ($(MAKECMDGOALS),clean)
+ifneq ($(strip $(C_DEPS)),)
+-include $(C_DEPS)
+endif
+endif
+# Add inputs and outputs from these tool invocations to the build variables 
+# All Target
+all: oclh_builder
+# Tool invocations
+oclh_builder: DEFS= -D__OCLH_BUILDER_FLAG -D__OCLH_BUILD_LOG_TO_STDOUT_FLAG
+oclh_builder: OUTPUTNAME=$(OCLH_BUILDER_NAME)
+oclh_builder: $(OBJS) $(USER_OBJS)
+	@echo 'Building target: $@'
+	@echo 'Invoking C Linker'
+	$(LINKER) $(LDFLAGS) $(MISCFLAGS) $(LIBS) \
+	-o "$(TARGETDIR)/$(OUTPUTNAME)" $(OBJS) $(USER_OBJS)
+	@echo 'Finished building target: $@'
+	@echo ' '
+# Other Targets
+clean:
+	$(RM) $(OBJS)$(C_DEPS) \
+	$(TARGETDIR)/$(OCLH_BUILDER_NAME) \
+	$(TARGETDIR)/$(OCLH_COMPILER_NAME) \
+	$(TARGETDIR)/$(OCLH_LINKER_NAME) \
+	$(TARGETDIR)/$(OCLH_LIBRARY_NAME)
+	@echo ' '
+
+.PHONY: all clean dependents
+
+-include makefile.targets

examples/simple_kernel/simple_kernel.clc

@@ -0,0 +1,7 @@
+__kernel void krSimpleKernel(__global  float* pA,
+                             __global  float* pB,
+                             __private long unsigned int size) {
+    long unsigned int idx=get_global_id(0);
+    if(idx<size) pB[idx]=pA[idx]*pA[idx];
+    return;
+}

src/include_d/oclh_d_mem_alloc.clh

@@ -0,0 +1,314 @@
+/*
+ * oclh_d_mem_alloc.clh
+ *      Author: havock
+ */
+#ifndef OCLH_D_MEM_ALLOC_CLH_
+#define OCLH_D_MEM_ALLOC_CLH_ 1
+#include <oclh_hd_std_types.clh>
+#ifdef _OCLH_OCL_COMPILER_
+/*
+ * common heap macros
+ */
+#ifndef _GDM__PRIVATE_MEM_HEAP_SZ_BYTES
+/* size of private memory heap size
+ * accessible from inside of a kernel only */
+#define _GDM__PRIVATE_MEM_HEAP_SZ_BYTES         65536
+#endif
+#ifndef _GDM__GLOBAL_EXCL_MEM_HEAP_SZ_BYTES
+/* size of global memory heap size
+ * accessible from anywhere */
+#define _GDM__GLOBAL_EXCL_MEM_HEAP_SZ_BYTES     65536
+#endif
+#ifndef _GDM__MEM_ALIGN_BYTES
+/* align
+ * some architectures returns memory errors
+ * if pointer is not aligned */
+#define _GDM__MEM_ALIGN_BYTES                   8
+#endif
+/* errors */
+#define _GDM__PRIVATE_HEAP_ALLOC_ERR            127
+#define _GDM__GLOBAL_HEAP_ALLOC_ERR             128
+
+/*
+ * technical heap macros
+ */
+#define __GDM__HEAP_HEADER_SZ_BYTES 32
+#define __GDM__MHI_SIZE 0
+#define __GDM__MHI_PREV 1
+#define __GDM__MHI_CURR 2
+#define __GDM__MHI_NEXT 3
+#define \
+    __GDM_STRICTLY_TECHNICAL_common_malloc_MACROS(_Region_, _Type_, \
+                                                  _Ptr_, _Size_, _Align_, \
+                                                  _Heap_u8Ptr_, _Heap_Sz_B_, \
+                                                  _Hdr_Sz_B_) \
+{ \
+    union { _Region_ uint8_t* u8Ptr; \
+            _Region_ uint64_t* u64Ptr;} LastMH; \
+    __private uint64_t u64LastSz,u64LastPr,u64LastCr; \
+    __private uint64_t u64LastNx=0ul; \
+    __private uint64_t u64HeapTreshold=(_Heap_Sz_B_)/(_Hdr_Sz_B_); \
+    __private uint64_t u64RealSz=(_Size_); \
+    if(u64RealSz%(_Align_)) u64RealSz+=((_Align_)-(u64RealSz%(_Align_))); \
+    do { \
+        LastMH.u8Ptr=&((_Heap_u8Ptr_)[u64LastNx]); \
+        u64LastSz=LastMH.u64Ptr[__GDM__MHI_SIZE]; \
+        u64LastPr=LastMH.u64Ptr[__GDM__MHI_PREV]; \
+        u64LastCr=LastMH.u64Ptr[__GDM__MHI_CURR]; \
+        u64LastNx=LastMH.u64Ptr[__GDM__MHI_NEXT]; \
+        u64HeapTreshold--; \
+    } while( \
+            u64LastNx && \
+            (u64LastNx-(u64LastCr+(_Hdr_Sz_B_)+u64LastSz)) < \
+                ((_Hdr_Sz_B_)+u64RealSz) && \
+            u64HeapTreshold && \
+            u64LastNx<(_Heap_Sz_B_) \
+           ); \
+    if( \
+       !u64HeapTreshold || \
+       u64LastNx>=(_Heap_Sz_B_) || \
+       (!u64LastNx && ((_Hdr_Sz_B_)+u64RealSz) > \
+           ((_Heap_Sz_B_)-(u64LastCr+(_Hdr_Sz_B_)+u64LastSz))) \
+      ) { \
+        union { uint64_t u64Zero; \
+                _Region_ _Type_ * typePtr; } Memory; \
+        Memory.u64Zero=0ul; \
+        (_Ptr_)=Memory.typePtr; \
+    } else { \
+        union { _Region_ uint8_t* u8Ptr; \
+                _Region_ uint64_t* u64Ptr; } NewMH; \
+        __private uint64_t u64NewCr=u64LastCr+(_Hdr_Sz_B_)+u64LastSz; \
+        NewMH.u8Ptr=&((_Heap_u8Ptr_)[u64NewCr]); \
+        NewMH.u64Ptr[__GDM__MHI_SIZE] =u64RealSz; \
+        NewMH.u64Ptr[__GDM__MHI_PREV] =u64LastCr; \
+        NewMH.u64Ptr[__GDM__MHI_CURR] =u64NewCr; \
+        NewMH.u64Ptr[__GDM__MHI_NEXT] =u64LastNx; \
+        LastMH.u64Ptr[__GDM__MHI_NEXT]=u64NewCr; \
+        if(u64LastNx) { \
+            union { _Region_ uint8_t* u8Ptr; \
+                    _Region_ uint64_t* u64Ptr; } \
+                NextMH; \
+            NextMH.u8Ptr=&((_Heap_u8Ptr_)[u64LastNx]); \
+            NextMH.u64Ptr[__GDM__MHI_PREV]=u64NewCr; \
+        } \
+        { \
+            union { _Region_ uint8_t* u8Ptr; \
+                    _Region_ _Type_ * typePtr; } \
+                Memory; \
+            Memory.u8Ptr=&(_Heap_u8Ptr_)[u64NewCr+(_Hdr_Sz_B_)]; \
+            (_Ptr_)=Memory.typePtr; \
+        } \
+    } \
+}
+#define \
+    __GDM_STRICTLY_TECHNICAL_common_free_MACROS(_Region_, _Type_, _Ptr_, \
+                                                _Heap_u8Ptr_,_Heap_Sz_B_, \
+                                                _Hdr_Sz_B_) \
+{ \
+    union { _Region_ uint8_t* u8Ptr; \
+            _Region_ uint64_t* u64Ptr; \
+            _Region_ _Type_ * typePtr; } Memory; \
+    Memory.typePtr=(_Ptr_); \
+    Memory.u8Ptr=&Memory.u8Ptr[-_Hdr_Sz_B_]; \
+    __private const uint64_t u64PrevOff=Memory.u64Ptr[__GDM__MHI_PREV]; \
+    __private const uint64_t u64NextOff=Memory.u64Ptr[__GDM__MHI_NEXT]; \
+    if(u64PrevOff!=(-1)) { \
+        if(u64NextOff) { \
+            union { _Region_ uint8_t* u8Ptr; \
+                    _Region_ uint64_t* u64Ptr; } NextMH; \
+            NextMH.u8Ptr=&((_Heap_u8Ptr_)[u64NextOff]); \
+            NextMH.u64Ptr[__GDM__MHI_PREV]=u64PrevOff; \
+        } \
+        { \
+            union { _Region_ uint8_t* u8Ptr; \
+                    _Region_ uint64_t* u64Ptr; } PrevMH; \
+            PrevMH.u8Ptr=&((_Heap_u8Ptr_)[u64PrevOff]); \
+            PrevMH.u64Ptr[__GDM__MHI_NEXT]=u64NextOff; \
+        } \
+    } \
+    { \
+        union { _Region_ uint64_t u64Zero; \
+                _Region_ _Type_ * typePtr; } Memory; \
+        Memory.u64Zero=0ul; \
+        (_Ptr_)=Memory.typePtr; \
+    } \
+}
+#define \
+    __GDM_STRICTLY_TECHNICAL_print_heap_state_MACROS(_Region_, _Heap_u8Ptr_, \
+                                                     _Heap_Sz_B_, _Hdr_Sz_B_) \
+{ \
+    __private const uint64_t u64Idx=get_global_id(0); \
+    union { _Region_ const uint8_t* u8Ptr; \
+            _Region_ const uint64_t* u64Ptr; } MemHeader; \
+    __private uint64_t u64Size,u64PrevOff,u64CurrOff; \
+    __private uint64_t u64NextOff=0ul; \
+    __private uint64_t i=0ul; \
+    do { \
+        MemHeader.u8Ptr=&_Heap_u8Ptr_[u64NextOff]; \
+        u64Size=MemHeader.u64Ptr[__GDM__MHI_SIZE]; \
+        u64PrevOff=MemHeader.u64Ptr[__GDM__MHI_PREV]; \
+        u64CurrOff=MemHeader.u64Ptr[__GDM__MHI_CURR]; \
+        u64NextOff=MemHeader.u64Ptr[__GDM__MHI_NEXT]; \
+        printf("Thread #%lu | NULL=%lu Segment #%lu " \
+               "[ size:%lu prev:%lu curr:%lu next:%lu actptr:%lu ]\n", \
+               u64Idx,(uint64_t) _Heap_u8Ptr_,i, \
+               u64Size,u64PrevOff,u64CurrOff,u64NextOff, \
+               (uint64_t) &_Heap_u8Ptr_[u64CurrOff+_Hdr_Sz_B_]); \
+        i++; \
+    } while (u64NextOff); \
+}
+#define __GDM_malloc_SUBST(__GDM_MALLOC_SUBST_MAC_Region) \
+                                           malloc##__GDM_MALLOC_SUBST_MAC_Region
+#define __GDM_free_SUBST(__GDM_FREE_SUBST_MAC_Region) \
+                                               free##__GDM_FREE_SUBST_MAC_Region
+#define __GDM_print_state_SUBST(__GDM_PRINT_STATE_SUBST_MAC_Region) \
+                           Print##__GDM_PRINT_STATE_SUBST_MAC_Region##_HeapState
+
+/*
+ * usable heap macros
+ */
+#define _GDM___private_heap_init() \
+    __private uint8_t \
+        __GD____private_Heap_u8Ptr[_GDM__PRIVATE_MEM_HEAP_SZ_BYTES]; \
+    { \
+        union { __private uint8_t* u8Ptr; \
+                __private uint64_t* u64Ptr; } __GD__MemHeader; \
+        __GD__MemHeader.u8Ptr=__GD____private_Heap_u8Ptr; \
+        __GD__MemHeader.u64Ptr[__GDM__MHI_SIZE]=0ul; \
+        __GD__MemHeader.u64Ptr[__GDM__MHI_PREV]=(-1); \
+        __GD__MemHeader.u64Ptr[__GDM__MHI_CURR]=0ul; \
+        __GD__MemHeader.u64Ptr[__GDM__MHI_NEXT]=0ul; \
+    }
+#define _GDM___global_heap_init(glblexclu8Ptr) \
+        __global uint8_t* const __GD____global_Heap_u8Ptr=glblexclu8Ptr; \
+        { \
+            union { __global uint8_t* u8Ptr; \
+                    __global uint64_t* u64Ptr; } __GD__MemHeader; \
+            __GD__MemHeader.u8Ptr=__GD____global_Heap_u8Ptr; \
+            __GD__MemHeader.u64Ptr[__GDM__MHI_SIZE]=0ul; \
+            __GD__MemHeader.u64Ptr[__GDM__MHI_PREV]=(-1); \
+            __GD__MemHeader.u64Ptr[__GDM__MHI_CURR]=0ul; \
+            __GD__MemHeader.u64Ptr[__GDM__MHI_NEXT]=0ul; \
+    }
+#define _GDM_heap_PROTO(__GDM_HEAP_PROTO_MAC_Region) \
+    __GDM_HEAP_PROTO_MAC_Region uint8_t* const \
+        __GD__##__GDM_HEAP_PROTO_MAC_Region##_Heap_u8Ptr
+#define _GDM_heap_ARG(__GDM_HEAP_ARG_MAC_Region) \
+    __GD__##__GDM_HEAP_ARG_MAC_Region##_Heap_u8Ptr
+#define _GDM_cast_pointer(__GDM_CAST_MACROS_Region, __GDM_CAST_MACROS_dstType, \
+                          __GDM_CAST_MACROS_dstPtr, __GDM_CAST_MACROS_srcType, \
+                          __GDM_CAST_MACROS_srcPtr) { \
+    union { \
+        __GDM_CAST_MACROS_Region __GDM_CAST_MACROS_dstType * dstPtr; \
+        __GDM_CAST_MACROS_Region __GDM_CAST_MACROS_srcType * srcPtr; \
+    } __GD_CAST_MACROS_CastUnion; \
+    __GD_CAST_MACROS_CastUnion.srcPtr=(__GDM_CAST_MACROS_srcPtr); \
+    (__GDM_CAST_MACROS_dstPtr)=__GD_CAST_MACROS_CastUnion.dstPtr; \
+}
+#define _GDM_malloc(__GDM_MALLOC_MAC_Region, __GDM_MALLOC_MAC_Type, \
+                    __GDM_MALLOC_MAC_Ptr, __GDM_MALLOC_MAC_Size) \
+{ \
+    _GDM_cast_pointer(__GDM_MALLOC_MAC_Region, \
+                      __GDM_MALLOC_MAC_Type, (__GDM_MALLOC_MAC_Ptr), uint8_t, \
+                      __GDM_malloc_SUBST(__GDM_MALLOC_MAC_Region) \
+                          ((__GD__##__GDM_MALLOC_MAC_Region##_Heap_u8Ptr), \
+                      (__GDM_MALLOC_MAC_Size))); \
+}
+
+#define _GDM_free(__GDM_MALLOC_MAC_Region, \
+                  __GDM_MALLOC_MAC_Type, __GDM_MALLOC_MAC_Ptr) \
+{ \
+    __GDM_MALLOC_MAC_Region uint8_t** _GD_free_MACROS_VAR_ppu8Ptr; \
+    _GDM_cast_pointer(__GDM_MALLOC_MAC_Region, uint8_t*, \
+                      _GD_free_MACROS_VAR_ppu8Ptr, __GDM_MALLOC_MAC_Type*, \
+                      &(__GDM_MALLOC_MAC_Ptr)); \
+    __GDM_free_SUBST(__GDM_MALLOC_MAC_Region) \
+        ((__GD__ ## __GDM_MALLOC_MAC_Region ## _Heap_u8Ptr), \
+        _GD_free_MACROS_VAR_ppu8Ptr); \
+}
+#define _GDM_print_heap_state(__GDM_PRINT_HEAP_STATE_MAC_Region) \
+{ \
+    __GDM_print_state_SUBST(__GDM_PRINT_HEAP_STATE_MAC_Region) \
+        ((__GD__ ## __GDM_PRINT_HEAP_STATE_MAC_Region ## _Heap_u8Ptr)); \
+}
+
+__private uint8_t* malloc__private(__private       uint8_t* const pu8HeapPtr,
+                                             const uint64_t       u64Size);
+__global  uint8_t* malloc__global( __global        uint8_t* const pu8HeapPtr,
+                                             const uint64_t       u64Size);
+void free__private(__private uint8_t* const pu8HeapPtr,
+                   __private uint8_t**      ppu8Ptr);
+void free__global( __global  uint8_t* const pu8HeapPtr,
+                   __global  uint8_t**      ppu8Ptr);
+int32_t Print__private_HeapState(__private const uint8_t* const pu8HeapPtr);
+int32_t Print__global_HeapState( __global  const uint8_t* const pu8HeapPtr);
+
+__private uint8_t* malloc__private(__private       uint8_t* const pu8HeapPtr,
+                                             const uint64_t       u64Size) {
+    __private uint8_t* pu8Res;
+    __GDM_STRICTLY_TECHNICAL_common_malloc_MACROS(__private, uint8_t,
+                                                  pu8Res, u64Size,
+                                                  _GDM__MEM_ALIGN_BYTES,
+                                                  pu8HeapPtr,
+                                                  _GDM__PRIVATE_MEM_HEAP_SZ_BYTES,
+                                                  __GDM__HEAP_HEADER_SZ_BYTES);
+    return(pu8Res);
+}
+__global uint8_t* malloc__global(__global       uint8_t* const pu8HeapPtr,
+                                          const uint64_t       u64Size) {
+    __global uint8_t* pu8Res;
+    __GDM_STRICTLY_TECHNICAL_common_malloc_MACROS(__global, uint8_t,
+                                                  pu8Res, u64Size,
+                                                  _GDM__MEM_ALIGN_BYTES,
+                                                  pu8HeapPtr,
+                                                  _GDM__GLOBAL_EXCL_MEM_HEAP_SZ_BYTES,
+                                                  __GDM__HEAP_HEADER_SZ_BYTES);
+    return(pu8Res);
+}
+void free__private(__private uint8_t* const pu8HeapPtr,
+                   __private uint8_t**      ppu8Ptr) {
+    if(*ppu8Ptr) {
+        __GDM_STRICTLY_TECHNICAL_common_free_MACROS(__private, uint8_t,
+                                                    *ppu8Ptr, pu8HeapPtr,
+                                                    _GDM__PRIVATE_MEM_HEAP_SZ_BYTES,
+                                                    __GDM__HEAP_HEADER_SZ_BYTES);
+    }
+    return;
+}
+void free__global(__global uint8_t* const pu8HeapPtr,
+                  __global uint8_t**      ppu8Ptr) {
+    if(*ppu8Ptr) {
+        __GDM_STRICTLY_TECHNICAL_common_free_MACROS(__global, uint8_t,
+                                                    *ppu8Ptr, pu8HeapPtr,
+                                                    _GDM__GLOBAL_EXCL_MEM_HEAP_SZ_BYTES,
+                                                    __GDM__HEAP_HEADER_SZ_BYTES);
+    }
+    return;
+}
+int32_t Print__private_HeapState(__private const uint8_t* const pu8HeapPtr) {
+    __GDM_STRICTLY_TECHNICAL_print_heap_state_MACROS(__private, pu8HeapPtr,
+                                                     _GDM__PRIVATE_MEM_HEAP_SZ_BYTES,
+                                                     __GDM__HEAP_HEADER_SZ_BYTES);
+    return(0);
+}
+int32_t Print__global_HeapState(__global const uint8_t* const pu8HeapPtr) {
+    __GDM_STRICTLY_TECHNICAL_print_heap_state_MACROS(__global, pu8HeapPtr,
+                                                     _GDM__PRIVATE_MEM_HEAP_SZ_BYTES,
+                                                     __GDM__HEAP_HEADER_SZ_BYTES);
+    return(0);
+}
+
+// flt32_t _gdf_atomicAdd_f32(__global volatile flt32_t *pfPtr, flt32_t fVal) {
+//     __global volatile uint32_t* pu32Ptr;
+//     union { uint32_t u32; flt32_t f32; } nxt,exp,cur;
+//     _GDM_cast_pointer(__global volatile ,uint32_t,pu32Ptr,flt32_t,pfPtr);
+//     cur.f32=*pfPtr;
+//     do {
+//         exp.f32=cur.f32; nxt.f32=exp.f32+fVal;
+//         cur.u32=atomic_cmpxchg(pu32Ptr,exp.u32,nxt.u32);
+//     } while(cur.u32!=exp.u32);
+//     return(cur.f32);
+// }
+
+#endif /* _OCLH_OCL_COMPILER_ */
+#endif /* OCLH_D_MEM_ALLOC_CLH_ */

src/include_d/oclh_d_srr.clh

@@ -0,0 +1,162 @@
+/*
+ * oclh_d_srr.clh
+ *      Author: havock
+ */
+#ifndef OCLH_D_SRR_CLH_DECLS_
+#define OCLH_D_SRR_CLH_DECLS_ 1
+#include <oclh_hd_std_types.clh>
+
+#if defined(_OCLH_OCL_HOST_HEADERS_) || defined(_OCLH_OCL_COMPILER_)
+/*
+ * current section is available for host just for memory size calculations
+ */
+#pragma pack(push,1)
+typedef struct _GDT_VECTOR_SET_DESCRIPTION {
+    __global flt32_t* pf32V;
+             uint64_t u64VDim;
+             uint64_t u64NofVs;
+} _GDT_VS_DESC;
+typedef struct _GDT_CONST_VECTOR_SET_DESCRIPTION {
+    __global const flt32_t* pf32V;
+                   uint64_t u64VDim;
+                   uint64_t u64NofVs;
+} _GDT_CVS_DESC;
+#pragma pack(pop)
+#endif /* defined(_OCLH_OCL_HOST_HEADERS_) || defined(_OCLH_OCL_COMPILER_) */
+
+#ifdef _OCLH_OCL_COMPILER_
+_GDT_VS_DESC _gdf_declVecSetFromTermOffset(__global flt32_t* const pf32V,
+                                           __private const uint64_t u64VDim,
+                                           __private const uint64_t u64VOffset,
+                                           __private const uint64_t u64NofVecs);
+_GDT_CVS_DESC _gdf_declConstVecSetFromTermOffset(
+                                     __global  const flt32_t* const pf32V,
+                                     __private const uint64_t       u64VDim,
+                                     __private const uint64_t       u64VOffset,
+                                     __private const uint64_t       u64NofVecs);
+int32_t _gdf_dropOutVecs_f32_g(__private const _GDT_VS_DESC VS);
+flt32_t _gdf_euclDst_f32_pp(__private const flt32_t* const pA,
+                            __private const flt32_t* const pB,
+                            __private const uint64_t       Dim);
+flt32_t _gdf_euclDst_f32_gg(__global  const flt32_t* const pA,
+                            __global  const flt32_t* const pB,
+                            __private const	uint64_t       Dim);
+int32_t _gdf_copyVec_f32_pp(__private       flt32_t* const pDst,
+                            __private const flt32_t* const pSrc,
+                            __private const uint64_t       Dim);
+int32_t _gdf_copyVec_f32_gg(__global        flt32_t* const pDst,
+                            __global  const flt32_t* const pSrc,
+                            __private const uint64_t       Dim);
+int32_t _gdf_copyVec_f32_pg(__private       flt32_t* const pDst,
+                            __global  const flt32_t* const pSrc,
+                            __private const uint64_t       Dim);
+uint32_t CRC32b(__global  const uint8_t* const pu8Data,
+                __private const uint64_t u64Sz);
+#endif /* _OCLH_OCL_COMPILER_ */
+
+#endif /* OCLH_D_SRR_CLH_DECLS_ */
+
+
+
+
+#ifdef _OCLH_OCL_COMPILER_
+#ifndef OCLH_D_SRR_CLH_IMPLS_
+#define OCLH_D_SRR_CLH_IMPLS_ 1
+
+_GDT_VS_DESC _gdf_declVecSetFromTermOffset(__global  flt32_t* const pf32V,
+                                           __private const uint64_t u64VDim,
+                                           __private const uint64_t u64VOffset,
+                                           __private const uint64_t u64NofVecs){
+    __private _GDT_VS_DESC desc={
+        .pf32V=&pf32V[u64VOffset*u64VDim],
+        .u64VDim=u64VDim,
+        .u64NofVs=u64NofVecs
+    };
+    return(desc);
+}
+_GDT_CVS_DESC _gdf_declConstVecSetFromTermOffset(
+                                           __global  const flt32_t* const pf32V,
+                                           __private const uint64_t u64VDim,
+                                           __private const uint64_t u64VOffset,
+                                           __private const uint64_t u64NofVecs){
+    __private _GDT_CVS_DESC desc={
+        .pf32V=&pf32V[u64VOffset*u64VDim],
+        .u64VDim=u64VDim,
+        .u64NofVs=u64NofVecs
+    };
+    return(desc);
+}
+int32_t _gdf_dropOutVecs_f32_g(__private const _GDT_VS_DESC VS) {
+    /*
+     * use this function for one thread only
+     * using the function in multithreaded case could lead to
+     * an inseparable mix of the output result
+     */
+    __private uint64_t i=0ul;
+    for(i=0ul; i<VS.u64NofVs; i++) {
+        __private uint64_t j=0ul;
+        __global const flt32_t* const pf32CurV=&VS.pf32V[i*VS.u64VDim];
+        printf("# %lu # ",i);
+        for(j=0; j<VS.u64VDim; j++) printf("%f ",pf32CurV[j]); printf("\n");
+    }
+    printf("\n");
+    return(0);
+}
+flt32_t _gdf_euclDst_f32_pp(__private const flt32_t* const pA,
+                            __private const flt32_t* const pB,
+                            __private const uint64_t       Dim) {
+    __private flt32_t S=0.0f, D=0.0f;
+    __private uint64_t i=0ul;
+    for(i=0ul; i<Dim; i++) { D=pA[i]-pB[i]; S+=(D*D); };
+    return(sqrt(S));
+}
+flt32_t _gdf_euclDst_f32_gg(__global  const flt32_t* const pfA,
+                            __global  const flt32_t* const pfB,
+                            __private const uint64_t       Dim) {
+    __private flt32_t S=0.0f, D=0.0f;
+    __private uint64_t i=0ul;
+    for(i=0; i<Dim; i++) { D=pfA[i]-pfB[i]; S+=(D*D); };
+    return(sqrt(S));
+}
+int32_t _gdf_copyVec_f32_pp(__private       flt32_t* const pDst,
+                            __private const flt32_t* const pSrc,
+                            __private const uint64_t       Dim) {
+    __private uint64_t i=0ul;
+    for(i=0ul; i<Dim; i++) pDst[i]=pSrc[i];
+    return(0);
+}
+int32_t _gdf_copyVec_f32_gg(__global        flt32_t* const pDst,
+                            __global  const flt32_t* const pSrc,
+                            __private const uint64_t       Dim) {
+    __private uint64_t i=0ul;
+    for(i=0ul; i<Dim; i++) pDst[i]=pSrc[i];
+    return(0);
+}
+int32_t _gdf_copyVec_f32_pg(__private       flt32_t* const pDst,
+                            __global  const flt32_t* const pSrc,
+                            __private const uint64_t       Dim) {
+    __private uint64_t i=0ul;
+    for(i=0;i<Dim;i++) pDst[i]=pSrc[i];
+    return(0);
+}
+
+uint32_t CRC32b(__global  const uint8_t* const pu8Data,
+                __private const uint64_t       u64Sz) {
+    __private uint64_t i=0ul;
+    __private uint32_t byte=0u,
+                       mask=0u;
+                       crc=0xFFFFFFFF;
+    for(i=0ul; i<u64Sz; i++) {
+        __private int32_t j=0;
+        byte=(uint32_t)(pu8Data[i]);
+        crc=crc^byte;
+        for(j=7; j>=0; j--) {
+            mask=-(crc & 1);
+            crc=(crc >> 1)^(0xEDB88320 & mask);
+        }
+    }
+    return(crc);
+}
+
+#endif /* OCLH_D_SRR_CLH_IMPLS_ */
+#endif /* _OCLH_OCL_COMPILER_ */

src/include_h/oclh.h

@@ -0,0 +1,28 @@
+/*
+ * oclh.h
+ *      Author: havock
+ */
+#ifndef OCLH_H_
+#define OCLH_H_ 1
+
+/*
+ * Соглашения:
+ * _ghf_*()  - gpu host function - функция только для хоста
+ * _gdf_*()  - gpu device function - функция только для устройства
+ * _ghdf_*() - gpu host-device function - функция единая для хоста и устройства
+ * _GHM_*    - gpu host macro - макрос только для хоста
+ * _GDM_*    - gpu device macro - макрос только для устройства
+ * _GHT_*    - gpu host type
+ * _GHT_*    - gpu device type
+ * _GHE_*    - gpu host enum
+ *
+ * _sqf_*()  - gpu device squad function
+ * _SQT_*    - gpu device squad type
+ * _SQCMD_*  - gpu device squad command macros
+ */
+#include <oclh_h_settings.h>
+#include <oclh_h_ws_base.h>
+#include <oclh_h_internals.h>
+#include <oclh_h_externals.h>
+
+#endif /* OCLH_H_ */

src/include_h/oclh_cc.h

@@ -0,0 +1,11 @@
+/*
+ * oclh_cc.h
+ *
+ *  Created on: Jun 1, 2015
+ *      Author: havock
+ */
+
+#ifndef OCLH_CC_H_
+#define OCLH_CC_H_ 1
+
+#endif /* OCLH_CC_H_ */

src/include_h/oclh_h_base_clapi_strings.h

@@ -0,0 +1,19 @@
+/*
+ * oclh_h_base_clapi_wrappers.h
+ *      Author: havock
+ */
+#ifndef OCLH_H_BASE_CLAPI_WRAPPERS_H_
+#define OCLH_H_BASE_CLAPI_WRAPPERS_H_ 1
+#include <CL/opencl.h>
+
+typedef enum _GHE_REC_TYPE {
+    _GHE_RT_POSTFIX = 0,
+    _GHE_RT_PREFIX =  1
+} _GHE_RECTYPE;
+
+char* _ghf_CLAPIErrString(      const cl_int         clErr);
+char* _ghf_PlatfPropString(     const cl_uint        clProp);
+char* _ghf_DevPropSuffixString( const cl_device_info clDevInf,
+                                const _GHE_RECTYPE   rt);
+
+#endif /* OCLH_H_BASE_CLAPI_WRAPPERS_H_ */

src/include_h/oclh_h_base_defs.h

@@ -0,0 +1,13 @@
+/*
+ * oclh_h_base_defs.h
+ *      Author: havock
+ */
+
+#ifndef OCLH_H_BASE_DEFS_H_
+#define OCLH_H_BASE_DEFS_H_ 1
+
+#define _GHM_UNDEFUINTVAL   ((uint64_t)(-1)) /* very controversial assumption */
+#define _GHM_UNDEFPTR           NULL
+#define _GHM_UNDEFBUILDSTATUS   1
+
+#endif /* OCLH_H_BASE_DEFS_H_ */

src/include_h/oclh_h_base_dev_clapi_wrappers.h

@@ -0,0 +1,29 @@
+/*
+ * oclh_h_base_dev_clapi_wrappers.h
+ *      Author: havock
+ */
+
+#ifndef OCLH_H_BASE_DEV_CLAPI_WRAPPERS_H_
+#define OCLH_H_BASE_DEV_CLAPI_WRAPPERS_H_ 1
+#include <CL/opencl.h>
+
+cl_bool   _ghf_getDevInf_bool(   const cl_device_id   clDev,
+                                 const cl_device_info InfoVal,
+                                       cl_int*  const pclErr);
+char*     _ghf_getDevInf_charptr(const cl_device_id   clDev,
+                                 const cl_device_info InfoVal,
+                                       cl_int*  const pclErr);
+cl_uint   _ghf_getDevInf_cluint( const cl_device_id   clDev,
+                                 const cl_device_info InfoVal,
+                                       cl_int*  const pclErr);
+cl_ulong  _ghf_getDevInf_clulong(const cl_device_id   clDev,
+                                 const cl_device_info InfoVal,
+                                       cl_int*  const pclErr);
+size_t    _ghf_getDevInf_size(   const cl_device_id   clDev,
+                                 const cl_device_info InfoVal,
+                                       cl_int*  const pclErr);
+uintptr_t _ghf_getDevInf_uintptr(const cl_device_id   clDev,
+                                 const cl_device_info InfoVal,
+                                       cl_int*  const pclErr);
+
+#endif /* OCLH_H_BASE_DEV_CLAPI_WRAPPERS_H_ */

src/include_h/oclh_h_base_log.h

@@ -0,0 +1,33 @@
+/*
+ * oclh_h_base_log.h
+ *
+ *  Created on: 17 March 2020
+ *      Author: havock
+ */
+
+#ifndef OCLH_H_BASE_LOG_H_
+#define OCLH_H_BASE_LOG_H_ 1
+#include <stdio.h>
+#include <stdint.h>
+#include <pthread.h>
+#include <oclh_h_errors.h>
+
+typedef enum _GHE_LOG_LEVEL {
+    _GHE_LOG_DEFAULT  = 0,
+    _GHE_LOG_KEY_PARS = 1,
+    _GHE_LOG_ALL      = 2
+} _GHE_LOGLVL;
+
+#pragma pack(push,1)
+typedef struct _GHT_LOGFILE {
+    FILE*            pfOut;
+    pthread_mutex_t* pMtx;
+} _GHT_LOG;
+#pragma pack(pop)
+
+_GHT_LOG _ghf_declLog(void);
+int32_t  _ghf_genrLog(_GHT_LOG* const pLog, const char* const pcFileName);
+int32_t  _ghf_isLogValid(const _GHT_LOG Log);
+int32_t  _ghf_wipeLog(_GHT_LOG* const pLog);
+
+#endif /* OCLH_H_BASE_LOG_H_ */

src/include_h/oclh_h_basetype_templates.h

@@ -0,0 +1,156 @@
+/*
+ * oclh_h_basetype_templates.h
+ *      Author: havock
+ */
+#ifndef OCLH_H_BASETYPE_TEMPLATES_H_
+#define OCLH_H_BASETYPE_TEMPLATES_H_ 1
+
+#define _GHM_BASETYPE_HD_DECLARATION(__GHM_BASETYPE_ARG_DEC__)                 \
+typedef struct _GHT_HOST_DEV_##__GHM_BASETYPE_ARG_DEC__##_PTR {                \
+    __GHM_BASETYPE_ARG_DEC__ * hPtr;                                           \
+    cl_mem   dPtr;                                                             \
+    uint64_t u64NofVals;                                                       \
+} _GHT_HD##__GHM_BASETYPE_ARG_DEC__ ;                                          \
+_GHT_HD##__GHM_BASETYPE_ARG_DEC__ _ghf_declHD_##__GHM_BASETYPE_ARG_DEC__(void);\
+int32_t _ghf_wdcAllocHD_##__GHM_BASETYPE_ARG_DEC__ (_GHT_WRKSET wSet,          \
+            _GHT_HD##__GHM_BASETYPE_ARG_DEC__ * const pBaseType,               \
+            const uint64_t u64NofVals);                                        \
+int32_t _ghf_wdcReAllocHD_##__GHM_BASETYPE_ARG_DEC__(_GHT_WRKSET wSet,         \
+            const _GHE_SYNC_TYPE SyncType,                                     \
+            _GHT_HD##__GHM_BASETYPE_ARG_DEC__ * const pBaseType,               \
+            const uint64_t u64NofVals);                                        \
+int32_t _ghf_wdcSyncHD_##__GHM_BASETYPE_ARG_DEC__(_GHT_WRKSET wSet,            \
+            const _GHE_SYNC_TYPE SyncType,                                     \
+            _GHT_HD##__GHM_BASETYPE_ARG_DEC__ * const pBaseType);              \
+int32_t _ghf_wdcFillHD_##__GHM_BASETYPE_ARG_DEC__(_GHT_WRKSET wSet,            \
+            _GHT_HD##__GHM_BASETYPE_ARG_DEC__ * const pBaseType,               \
+            const __GHM_BASETYPE_ARG_DEC__ Val);                               \
+int32_t _ghf_isHD_##__GHM_BASETYPE_ARG_DEC__##_Valid(                          \
+            const _GHT_HD##__GHM_BASETYPE_ARG_DEC__ BaseType);                 \
+int32_t _ghf_wipeHD_##__GHM_BASETYPE_ARG_DEC__(_GHT_WRKSET wSet,               \
+            _GHT_HD##__GHM_BASETYPE_ARG_DEC__ * const pBaseType);              \
+
+#define _GHM_BASETYPE_HD_IMPLEMENTATION(__GHM_BASETYPE_ARG_IMP__)              \
+_GHT_HD##__GHM_BASETYPE_ARG_IMP__ _ghf_declHD_##__GHM_BASETYPE_ARG_IMP__(void){\
+    _GHT_HD##__GHM_BASETYPE_ARG_IMP__ BaseType={ .hPtr=NULL,                   \
+                                                 .dPtr=NULL,.u64NofVals=0 };   \
+    return(BaseType);                                                          \
+}                                                                              \
+int32_t _ghf_wdcAllocHD_##__GHM_BASETYPE_ARG_IMP__(_GHT_WRKSET wSet,           \
+            _GHT_HD##__GHM_BASETYPE_ARG_IMP__ * const pBaseType,               \
+            const uint64_t u64NofVals) {                                       \
+    if(!pBaseType) return(_GHM_NULL_POINTER_RECEIVED_ERROR);                   \
+    _ghf_wipeHD_##__GHM_BASETYPE_ARG_IMP__ (wSet,pBaseType);                   \
+    pBaseType->hPtr=_ghf_wdcAllocHostBuf(wSet,                                 \
+                                  sizeof(__GHM_BASETYPE_ARG_IMP__)*u64NofVals);\
+    if(!pBaseType->hPtr) {                                                     \
+        _ghf_wipeHD_##__GHM_BASETYPE_ARG_IMP__ (wSet,pBaseType);               \
+        return(_GHM_HOST_MEMALLOC_ERROR);                                      \
+    }                                                                          \
+    pBaseType->dPtr=                                                           \
+        _ghf_wdcAllocDevBuf(wSet,CL_MEM_READ_WRITE,                            \
+            sizeof(__GHM_BASETYPE_ARG_IMP__)*u64NofVals,NULL);                 \
+    if(!pBaseType->dPtr) {                                                     \
+        _ghf_wipeHD_##__GHM_BASETYPE_ARG_IMP__ (wSet,pBaseType);               \
+        return(_GHM_DEVICE_MEMALLOC_ERROR);                                    \
+    }                                                                          \
+    pBaseType->u64NofVals=u64NofVals;                                          \
+    return(_GHM_OK);                                                           \
+}                                                                              \
+int32_t _ghf_wdcReAllocHD_##__GHM_BASETYPE_ARG_IMP__(_GHT_WRKSET wSet,         \
+            const _GHE_SYNC_TYPE SyncType,                                     \
+            _GHT_HD##__GHM_BASETYPE_ARG_IMP__ * const pBaseType,               \
+            const uint64_t u64NofVals) {                                       \
+    __GHM_BASETYPE_ARG_IMP__ * hTmpPtr=NULL;                                   \
+    if(!pBaseType) return(_GHM_NULL_POINTER_RECEIVED_ERROR);                   \
+    hTmpPtr=realloc(pBaseType->hPtr,                                           \
+                    sizeof(__GHM_BASETYPE_ARG_IMP__)*u64NofVals);              \
+    if(!hTmpPtr) {                                                             \
+        if(wSet.pvDat) wSet.pfnDatCleaner(wSet);                               \
+        _ghf_wipeHD_##__GHM_BASETYPE_ARG_IMP__ (wSet,pBaseType);               \
+        return(_GHM_HOST_MEMALLOC_ERROR);                                      \
+    }                                                                          \
+    pBaseType->hPtr=hTmpPtr;                                                   \
+    _ghf_freeDevZ(wSet, &pBaseType->dPtr);                                     \
+    pBaseType->dPtr=                                                           \
+        _ghf_wdcAllocDevBuf(wSet,CL_MEM_READ_WRITE,                            \
+                            sizeof(__GHM_BASETYPE_ARG_IMP__)*u64NofVals,NULL); \
+    if(!pBaseType->dPtr) {                                                     \
+        _ghf_wipeHD_##__GHM_BASETYPE_ARG_IMP__ (wSet,pBaseType);               \
+        return(_GHM_DEVICE_MEMALLOC_ERROR);                                    \
+    }                                                                          \
+    pBaseType->u64NofVals=u64NofVals;                                          \
+    { int32_t err=0;                                                           \
+      if((err= _ghf_wdcSyncHD_##__GHM_BASETYPE_ARG_IMP__ (wSet,                \
+                                          SyncType,pBaseType))) return(err); } \
+    return(_GHM_OK);                                                           \
+}                                                                              \
+int32_t _ghf_wdcSyncHD_##__GHM_BASETYPE_ARG_IMP__(_GHT_WRKSET wSet,            \
+            const _GHE_SYNC_TYPE SyncType,                                     \
+            _GHT_HD##__GHM_BASETYPE_ARG_IMP__ * const pBaseType) {             \
+    int32_t err=0;                                                             \
+    if(!pBaseType) return(_GHM_NULL_POINTER_RECEIVED_ERROR);                   \
+    switch(SyncType) {                                                         \
+        case _GHE_SYNC_HOST_TO_DEV:                                            \
+            {                                                                  \
+                char pcLogMsg[_GHM_MAXLEN_OF_LOGMSG];                          \
+                snprintf(pcLogMsg, _GHM_MAXLEN_OF_LOGMSG,                      \
+                         "%s/clEnqueueWriteBuffer", __func__);                 \
+                wSet.APIErr=clEnqueueWriteBuffer(wSet.Queue,                   \
+                    pBaseType->dPtr, CL_TRUE, 0ul,                             \
+                    sizeof(__GHM_BASETYPE_ARG_IMP__)*pBaseType->u64NofVals,    \
+                    pBaseType->hPtr, 0u, NULL, NULL);                          \
+                if((err=_ghf_wdcChkWS_APIErr(wSet,pcLogMsg,_GHM_FL))) {        \
+                    _ghf_wipeHD_##__GHM_BASETYPE_ARG_IMP__ (wSet,pBaseType);   \
+                    return(err);                                               \
+                }                                                              \
+            }                                                                  \
+            break;                                                             \
+        case _GHE_SYNC_DEV_TO_HOST:                                            \
+            {                                                                  \
+                char pcLogMsg[_GHM_MAXLEN_OF_LOGMSG];                          \
+                snprintf(pcLogMsg, _GHM_MAXLEN_OF_LOGMSG,                      \
+                         "%s/clEnqueueReadBuffer", __func__);                  \
+                wSet.APIErr=clEnqueueReadBuffer(wSet.Queue,                    \
+                    pBaseType->dPtr, CL_TRUE, 0ul,                             \
+                    sizeof(__GHM_BASETYPE_ARG_IMP__)*pBaseType->u64NofVals,    \
+                    pBaseType->hPtr, 0u, NULL, NULL);                          \
+                if((err=_ghf_wdcChkWS_APIErr(wSet,pcLogMsg,_GHM_FL))) {        \
+                    _ghf_wipeHD_##__GHM_BASETYPE_ARG_IMP__ (wSet,pBaseType);   \
+                    return(err);                                               \
+                }                                                              \
+            }                                                                  \
+            break;                                                             \
+        default:                                                               \
+            return(_GHM_UNKNOWN_SYNC_TYPE_ERROR);                              \
+            break;                                                             \
+    }                                                                          \
+    return(_GHM_OK);                                                           \
+}                                                                              \
+int32_t _ghf_wdcFillHD_##__GHM_BASETYPE_ARG_IMP__ (_GHT_WRKSET wSet,           \
+            _GHT_HD##__GHM_BASETYPE_ARG_IMP__ * const pBaseType,               \
+            const __GHM_BASETYPE_ARG_IMP__ Val) {                              \
+    int32_t err=0;                                                             \
+    uint64_t i=0ul;                                                            \
+    if(!pBaseType) return(_GHM_NULL_POINTER_RECEIVED_ERROR);                   \
+    for(i=0;i<pBaseType->u64NofVals;i++) pBaseType->hPtr[i]=Val;               \
+    if((err= _ghf_wdcSyncHD_##__GHM_BASETYPE_ARG_IMP__ (wSet,                  \
+        _GHE_SYNC_HOST_TO_DEV,pBaseType))) return(err);                        \
+    return(_GHM_OK);                                                           \
+}                                                                              \
+int32_t _ghf_isHD_##__GHM_BASETYPE_ARG_IMP__##_Valid(                          \
+                           const _GHT_HD##__GHM_BASETYPE_ARG_IMP__ BaseType) { \
+    return(BaseType.hPtr && BaseType.dPtr && BaseType.u64NofVals);             \
+}                                                                              \
+int32_t _ghf_wipeHD_##__GHM_BASETYPE_ARG_IMP__ (_GHT_WRKSET wSet,              \
+                        _GHT_HD##__GHM_BASETYPE_ARG_IMP__ * const pBaseType) { \
+    if(pBaseType) {                                                            \
+        _ghf_freeDevZ(wSet, &pBaseType->dPtr);                                 \
+        _ghf_freeHostZ(&pBaseType->hPtr);                                      \
+        pBaseType->u64NofVals=0ul;                                             \
+    }                                                                          \
+    return(_GHM_OK);                                                           \
+}                                                                              \
+
+#endif /* OCLH_H_BASETYPE_TEMPLATES_H_ */
+

src/include_h/oclh_h_clapi_callbacks.h

@@ -0,0 +1,15 @@
+/*
+ * oclh_clapi_callbacks.h
+ *      Author: havock
+ */
+#ifndef OCLH_H_CLAPI_CALLBACKS_H_
+#define OCLH_H_CLAPI_CALLBACKS_H_ 1
+#include <stdint.h>
+#include <CL/opencl.h>
+
+void CL_CALLBACK _ghf_CtxEvent(const char*  pcErrInfo,
+                               const void*  pvPrivInfo,
+                                     size_t szPrivInfo,
+                                     void*  pvUserData);
+
+#endif /* OCLH_H_CLAPI_CALLBACKS_H_ */

src/include_h/oclh_h_errors.h

@@ -0,0 +1,34 @@
+/*
+ * oclh_error_codes.h
+ *      Author: havock
+ */
+
+#ifndef OCLH_H_ERROR_CODES_H_
+#define OCLH_H_ERROR_CODES_H_ 1
+
+#define _GHM_OK                                     0
+#define _GHM_HOST_MEMALLOC_ERROR                 -150
+#define _GHM_DEVICE_MEMALLOC_ERROR               -151
+#define _GHM_UNKNOWN_SYNC_TYPE_ERROR             -152
+#define _GHM_OPEN_FILE_ERROR                     -153
+#define _GHM_SEEK_FILE_ERROR                     -154
+#define _GHM_READ_FILE_ERROR                     -155
+#define _GHM_WRITE_FILE_ERROR                    -156
+#define _GHM_EMPTY_STRING_RECEIVED_ERROR         -157
+#define _GHM_DECONSTRUCTION_OF_WORKSET_ERROR     -158
+#define _GHM_COMMAND_QUEUE_FLUSH_ERROR           -159
+#define _GHM_COMMAND_QUEUE_FINISH_ERROR          -160
+#define _GHM_UNAVALABLE_KERNEL_ERROR             -161
+#define _GHM_NO_PLATFORMS_AVAILABLE              -162
+#define _GHM_NO_DEVICES_AVAILABLE                -163
+#define _GHM_CREATING_KERNELS_ERROR              -164
+#define _GHM_INCONSISTED_LOG_ERROR               -165
+#define _GHM_NULL_POINTER_RECEIVED_ERROR         -166
+#define _GHM_ZERO_SIZE_RECEIVED_ERROR            -167
+#define _GHM_UNKNOWN_PROGRAM_SOURCE_ERROR        -168
+#define _GHM_INVALID_MAXSIZE_OF_WORK_GROUP_ERROR -169
+#define _GHM_INVALID_NUMBER_OF_CMP_UNITS_ERROR   -170
+
+#define _GHM_IMPOSSIBLE_SITUATION_ERROR      4030  /* the first weird four-digit
+                                                      (decimal) number */
+#endif /* OCLH_H_ERROR_CODES_H_ */

src/include_h/oclh_h_externals.h

@@ -0,0 +1,26 @@
+/*
+ * oclh_h_externals.h
+ *      Author: havock
+ */
+#ifndef OCLH_H_EXTERNALS_H_
+#define OCLH_H_EXTERNALS_H_ 1
+#include <oclh_h_ws_base.h>
+
+#define _GHM_FL   1
+#define _GHM_NOFL 0
+
+void*   _ghf_wdcAllocHostBuf(_GHT_WRKSET wSet, const size_t sz);
+void    _ghf_freeHostZ(void* const ppPtr);
+cl_mem  _ghf_wdcAllocDevBuf( _GHT_WRKSET  wSet,
+                       const cl_mem_flags cmFlags,
+                       const size_t       sz,
+                             void*  const pvHostPtr);
+void    _ghf_freeDevZ(_GHT_WRKSET wSet, cl_mem* const pCLMem);
+int32_t _ghf_addCharPtrToCharPtrList(      char*** const pppcLst,
+                                     const char*   const pcStr);
+int32_t _ghf_addFileToCharPtrList(      char*** const pppcLst,
+                                  const char*   const pcFileName,
+                                  const uint64_t      u64Align);
+void    _ghf_wipeCharPtrList(char*** const pppcLst);
+
+#endif /* OCLH_H_EXTERNALS_H_ */

src/include_h/oclh_h_internals.h

@@ -0,0 +1,15 @@
+/*
+ * oclh_internals.h
+ *
+ *  Created on: 17 March 2020
+ *      Author: havock
+ */
+#ifndef OCLH_INTERNALS_H_
+#define OCLH_INTERNALS_H_ 1
+#include <oclh_hd_std_types.clh>
+
+uint64_t __ghf_removePreNPostSpacesFromCharPtr(     char* const pcStr);
+uint64_t __ghf_replaceSpacesWithUnderscoreInCharPtr(char* const pcStr);
+int32_t __ghf_concatHeapStrAndCharPtr(char** ppcDst, const char* const pcStr);
+
+#endif

src/include_h/oclh_h_settings.h

@@ -0,0 +1,21 @@
+/*
+ * oclh_settings.h
+ *
+ *  Created on: Apr 29, 2016
+ *      Author: havock
+ */
+#ifndef INC_OCLH_SETTINGS_H_
+#define INC_OCLH_SETTINGS_H_ 1
+
+#define _GHM_LOG_PATH                    "."
+#define _GHM_MAX_PATH_LENGTH             1024
+#define _GHM_DEBUG_N_PROFILE             0
+#define _GHM_FLUSH_LOGS                  1
+#define _GHM_MAXLEN_OF_LOGMSG            1024
+#define _GHM_SHOW_PROGRAM_LISTING_IN_LOG 0
+#define _GHM_OCLH_BUILDER_LOG_FILENAME   "oclh_br.log"
+#define _GHM_OCLH_COMPILER_LOG_FILENAME  "oclh_cr.log"
+#define _GHM_OCLH_LINKER_LOG_FILENAME    "oclh_lr.log"
+#define _GHM_OCLH_COMMON_BUILD_OPTIONS   "-D_OCLH_OCL_COMPILER_"
+
+#endif /* INC_OCLH_SETTINGS_H_ */

src/include_h/oclh_h_srr.h

@@ -0,0 +1,94 @@
+/*
+ * oclh_h_srr.h
+ *      Author: havock
+ */
+
+#ifndef OCLH_H_SRR_H_
+#define OCLH_H_SRR_H_ 1
+#include <oclh_hd_std_types.clh>
+#include <oclh_h_ws_base.h>
+#include <oclh_h_basetype_templates.h>
+
+#define _OCLH_OCL_HOST_HEADERS_ 1
+#include <oclh_hd_srr.clh>
+#undef _OCLH_OCL_HOST_HEADERS_
+
+typedef enum _GHE_SYNC_TYPE {
+    _GHE_NO_SYNC          =0,
+    _GHE_SYNC_HOST_TO_DEV =1,
+    _GHE_SYNC_DEV_TO_HOST =2
+} _GHE_SYNC_TYPE;
+
+typedef int8_t   i8;
+typedef uint8_t  u8;
+typedef int16_t  i16;
+typedef uint16_t u16;
+typedef int32_t  i32;
+typedef uint32_t u32;
+typedef int64_t  i64;
+typedef uint64_t u64;
+typedef flt32_t  f32;
+typedef flt64_t  f64;
+typedef flt128_t  f128;
+
+#pragma pack(push,1)
+_GHM_BASETYPE_HD_DECLARATION(i8)
+_GHM_BASETYPE_HD_DECLARATION(u8)
+_GHM_BASETYPE_HD_DECLARATION(i16)
+_GHM_BASETYPE_HD_DECLARATION(u16)
+_GHM_BASETYPE_HD_DECLARATION(i32)
+_GHM_BASETYPE_HD_DECLARATION(u32)
+_GHM_BASETYPE_HD_DECLARATION(i64)
+_GHM_BASETYPE_HD_DECLARATION(u64)
+_GHM_BASETYPE_HD_DECLARATION(f32)
+_GHM_BASETYPE_HD_DECLARATION(f64)
+_GHM_BASETYPE_HD_DECLARATION(f128)
+#pragma pack(pop)
+
+#pragma pack(push,1)
+typedef struct _GHT_HOST_DEV_VOID_PTR {
+	uint8_t* hPtr;
+	cl_mem   dPtr;
+	uint64_t u64NofBytes;
+} _GHT_HDvoid;
+#pragma pack(pop)
+_GHT_HDvoid _ghf_declHD_void(void);
+int32_t _ghf_wdcAllocHD_void(  _GHT_WRKSET wSet, _GHT_HDvoid* const pHdv,
+                               const uint64_t u64NofBytes);
+int32_t _ghf_wdcReAllocHD_void(_GHT_WRKSET wSet, const _GHE_SYNC_TYPE SyncType,
+                               _GHT_HDvoid* const pHdv,
+                               const uint64_t u64NofBytes);
+int32_t _ghf_wdcSyncHD_void(   _GHT_WRKSET wSet, const _GHE_SYNC_TYPE SyncType,
+                               _GHT_HDvoid* const pHdv);
+int32_t _ghf_wdcFillHD_void(   _GHT_WRKSET wSet, _GHT_HDvoid* const pHdv,
+                               const uint8_t u8Val);
+int32_t _ghf_wipeHD_void(      _GHT_WRKSET wSet, _GHT_HDvoid* const pHdv);
+
+#pragma pack(push,1)
+typedef struct _GHT_HOST_DEV_VECS_FLT32 {
+    _GHT_HDf32 hdf32V;
+    uint64_t   u64Dim;
+    uint64_t   u64SemDim;
+} _GHT_HDVECS_F32;
+#pragma pack(pop)
+_GHT_HDVECS_F32 _ghf_declHDV_f32(void);
+int32_t  _ghf_wdcAllocHDV_f32(_GHT_WRKSET wSet, _GHT_HDVECS_F32* const pVhdf32,
+                              uint64_t u64Dim, const uint64_t u64SemDim,
+                              const uint64_t u64NofVecs);
+int32_t  _ghf_wdcSyncHDV_f32( _GHT_WRKSET wSet, const _GHE_SYNC_TYPE SyncType,
+                              _GHT_HDVECS_F32* const pVhdf32);
+int32_t  _ghf_wdcFillHDV_f32( _GHT_WRKSET wSet, _GHT_HDVECS_F32* const pVhdf32,
+                              const flt32_t f32Val);
+int32_t  _ghf_isHDV_f32_Valid(const _GHT_HDVECS_F32 Vhdf32);
+uint64_t _ghf_getNofVecsHDV_f32(const _GHT_HDVECS_F32 Vhdf32);
+int32_t  _ghf_saveHDV_f32_ToTxt(const _GHT_HDVECS_F32 vfVecs,
+                                const char* const pcFileName,
+                                const int32_t i32UseSemDimFlag);
+int32_t  _ghf_saveHDV_f32_ToTxtByHD_u64_Idxs(const _GHT_HDVECS_F32 vfVecs,
+                                             const _GHT_HDu64 hdu64IV,
+                                             const char* const pcFileName,
+                                             const int32_t i32UseSemDimFlag);
+int32_t  _ghf_wipeHDV_f32(_GHT_WRKSET wSet, _GHT_HDVECS_F32* const pVhdf32);
+
+#endif /* OCLH_H_SRR_H_ */
+

src/include_h/oclh_h_ws_base.h

@@ -0,0 +1,131 @@
+/*
+ * oclh_h_ws_base.h
+ *      Author: havock
+ */
+#ifndef OCLH_H_WS_BASE_H_
+#define OCLH_H_WS_BASE_H_ 1
+#include <sys/types.h>
+#include <stdlib.h>
+#include <CL/opencl.h>
+#include <oclh_h_base_log.h>
+
+#define __GHM_U64STRIPTO2B(___GHMARG_U64STRIPTO2B) \
+    (((uint64_t)(___GHMARG_U64STRIPTO2B))&((uint64_t)0x000000000000ffff))
+
+typedef enum _GHE_SOURCE_TYPE {
+    _GHE_HL_LISTINGS       = 0,
+    _GHE_IR_LISTINGS       = 1,
+    _GHE_SEPARATED_OBJECTS = 2,
+    _GHE_LINKED_OBJECTS    = 3
+} _GHE_SRCTYPE;
+
+typedef enum _GHE_BUILD_LOG_MODE {
+    _GHE_NO_BUILD_LOG                 = 0,
+    _GHE_BUILD_LOG_IN_WRKSET_LOG      = 1,
+    _GHE_BUILD_LOG_IN_SEPARATED_FILES = 2
+} _GHE_BUILD_LOG_MODE;
+
+#pragma pack(push,1)
+typedef struct _GHT_WORKINGSET {
+    void*            pwSetAddr;    /* self */
+    cl_command_queue Queue;        /* key pointer */
+    cl_program       Program;
+    size_t           szNofCmpUnits;/* number of  computational units
+                                      (a.k.a. processor core) */
+    size_t           szNofRPThrds; /* number of real parallel threads */
+    size_t           szMaxWGsz;    /* max number of work-items in a work-group
+                                      it represents the number of virtual
+                                      threads executed on a computational unit
+                                      (a.k.a. processor core) */
+    size_t           szNofWIinWG;
+    size_t           szNofAllWI;
+    size_t*          pszMemBytes;
+    _GHT_LOG         Log;
+    int32_t          i32EigenLogFlag;
+    cl_int           APIErr;
+    cl_kernel*       pKernels;
+    int32_t          KerErr;
+    cl_mem           d_pKerErr;
+    void*            pvDat;
+    int32_t          (*pfnDatCleaner)(const struct _GHT_WORKINGSET);
+} _GHT_WRKSET;
+#pragma pack(pop)
+
+_GHT_WRKSET _ghf_declWS(void);
+int32_t _ghf_genrWS(      _GHT_WRKSET*  const pwSet,
+                    const _GHT_LOG            Log,
+                    const int32_t             i32ExclusiveLogFlag,
+                    const cl_device_id        clWrkDev,
+                    const void*         const pvCLProgramSources,
+                    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) */
+                                                            );
+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_flfnWS(_GHT_WRKSET wSet);
+
+cl_device_id    _ghf_getWS_Dev(        _GHT_WRKSET wSet);
+cl_context      _ghf_getWS_Ctx(        _GHT_WRKSET wSet);
+cl_build_status _ghf_getWS_BuildStatus(_GHT_WRKSET wSet);
+cl_kernel       _ghf_getWS_KernByName( _GHT_WRKSET wSet,const char* pcKernName);
+cl_uint         _ghf_getWS_RefCntOfMem(_GHT_WRKSET wSet,const cl_mem clMem);
+size_t          _ghf_getWS_MaxWIinWG(  _GHT_WRKSET wSet);
+cl_uint         _ghf_getWS_MaxCmpUnits(_GHT_WRKSET wSet);
+int32_t         _ghf_saveWS_ProgramBinaries(_GHT_WRKSET wSet,
+                                            char* const pcOutputPrefix);
+int32_t         _ghf_wdcChkWS_APIErr (  _GHT_WRKSET wSet,
+                                  const char* const pcAPICall,
+                                  const int32_t i32FlashFlag);
+int32_t         _ghf_wdcChkWS_KerErr(   _GHT_WRKSET wSet,
+                                  const cl_kernel   clKer);
+int32_t         _ghf_wdcSetWS_KerErrToZero(_GHT_WRKSET wSet);
+#pragma pack(push,1)
+typedef struct _GHT_ALL_OF_WORKING_SETS {
+    _GHT_WRKSET* pWSet;
+    uint64_t u64NofWSs;
+    _GHT_LOG Log;
+    const void* const pvCLProgramSource;
+    const _GHE_SRCTYPE SourceType;
+    const _GHE_LOGLVL LogLevel;
+    const _GHE_BUILD_LOG_MODE BuildLogMode;
+} _GHT_AWSS;
+#pragma pack(pop)
+
+_GHT_AWSS _ghf_declAWSs(const _GHE_LOGLVL  LogLvl,
+                        const _GHE_BUILD_LOG_MODE BuildLogMode);
+int32_t _ghf_genrAWSs(      _GHT_AWSS* const pAWSs,
+                      const _GHT_LOG         Log,
+                      const void*      const pvCLProgramSources,
+                      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) */
+                                                           );
+int32_t _ghf_wipeAWSs(_GHT_AWSS* const pAWSs);
+
+int32_t _ghf_buildDevList(_GHT_WRKSET wSet, cl_device_id** ppDevLst);
+int32_t _ghf_wipeDevList(cl_device_id** ppDevLst);
+
+
+
+
+
+
+int32_t __ghf_setWS_TextProgramId(_GHT_WRKSET wSet,
+                                  char* const pcDst,
+                                  char* const pcOutputPrefix);
+
+#endif /* OCLH_H_WS_BASE_H_ */

src/include_h/oclh_h_ws_base_log.h

@@ -0,0 +1,56 @@
+/*
+ * oclh_h_wrkset_log.h
+ *      Author: havock
+ */
+#ifndef OCLH_H_WS_LOG_H_
+#define OCLH_H_WS_LOG_H_ 1
+#include <stdint.h>
+#include <oclh_h_ws_base.h>
+
+int32_t __ghf_logWS_MsgIgnoringLock(
+                           const _GHT_WRKSET wSet, const char* const pcLogMsg);
+int32_t _ghf_logWS_Msg(    const _GHT_WRKSET wSet, const char* const pcLogMsg);
+int32_t _ghf_logWS_Hdr(    const _GHT_WRKSET wSet, const char* const pcLogHdr);
+int32_t _ghf_logWS_Delim(  const _GHT_WRKSET wSet);
+int32_t _ghf_logWS_Raw(    const _GHT_WRKSET wSet, const char* const pcLogMsg);
+int32_t _ghf_logWS_APIErr( const _GHT_WRKSET wSet, const char* const pcAPICall);
+int32_t _ghf_logWS_UsedMem(const _GHT_WRKSET wSet);
+
+int32_t _ghf_logWS_DevInf_bool(  _GHT_WRKSET    wSet,
+                           const cl_device_id   clDev,
+                           const cl_device_info InfoVal,
+                           const char*    const pcPrefix,
+                           const char*    const pcName,
+                           const char*    const pcTrue,
+                           const char*    const pcFalse);
+int32_t _ghf_logWS_DevInf_charptr(_GHT_WRKSET   wSet,
+                           const cl_device_id   clDev,
+                           const cl_device_info InfoVal,
+                           const char*    const pcPrefix,
+                           const char*    const pcName,
+                           const char*    const pcPostfix);
+int32_t _ghf_logWS_DevInf_cluint(_GHT_WRKSET    wSet,
+                           const cl_device_id   clDev,
+                           const cl_device_info InfoVal,
+                           const char*    const pcPrefix,
+                           const char*    const pcName,
+                           const char*    const pcPostfix);
+int32_t _ghf_logWS_DevInf_clulong(_GHT_WRKSET   wSet,
+                           const cl_device_id   clDev,
+                           const cl_device_info InfoVal,
+                           const char*    const pcDevPrefix,
+                           const char*    const pcName,
+                           const char*    const pcPostfix);
+int32_t _ghf_logWS_DevInf_clulong_as_hex(_GHT_WRKSET wSet,
+                           const cl_device_id   clDev,
+                           const cl_device_info InfoVal,
+                           const char*    const pcPrefix,
+                           const char*    const pcName,
+                           const char*    const pcPostfix);
+int32_t _ghf_logWS_DevInf_size(  _GHT_WRKSET    wSet,
+                           const cl_device_id   clDev,
+                           const cl_device_info InfoVal,
+                           const char*    const pcDevPrefix,
+                           const char*    const pcName,
+                           const char*    const pcPostfix);
+#endif /* OCLH_H_WS_LOG_H_ */

src/include_h/oclh_h_ws_log_clapi_reps.h

@@ -0,0 +1,29 @@
+/*
+ * oclh_h_ws_log_clapi_reports.h
+ *      Author: havock
+ */
+#ifndef OCLH_H_WS_LOG_CLAPI_REPORTS_H_
+#define OCLH_H_WS_LOG_CLAPI_REPORTS_H_ 1
+#include <oclh_h_ws_base.h>
+
+int32_t _ghf_logWS_PlatfInfo(   const _GHT_WRKSET    wSet,
+                                const cl_platform_id Platform,
+                                const char*    const pcPrefix);
+int32_t _ghf_logWS_DevInfo(           _GHT_WRKSET    wSet,
+                                const cl_device_id   clDev,
+                                const _GHE_LOGLVL    LogLevel);
+int32_t _ghf_logWS_DevInfoShort(      _GHT_WRKSET    wSet,
+                                const cl_device_id   clDev,
+                                const char*    const pcPrefix);
+int32_t _ghf_logWS_ContextInfo(       _GHT_WRKSET    wSet,
+                                const cl_context     clCtx,
+                                const _GHE_LOGLVL    LogLevel);
+int32_t _ghf_logWS_BuildInfo(         _GHT_WRKSET    wSet,
+                                const cl_program     clProgram,
+                                const cl_device_id   clDev,
+                                const _GHE_LOGLVL    LogLevel,
+                                const _GHE_BUILD_LOG_MODE BuildLogMode);
+int32_t _ghf_logWS_KerInfo(           _GHT_WRKSET    wSet,
+                                const cl_kernel      clKer);
+
+#endif /* OCLH_H_WS_LOG_CLAPI_REPORTS_H_ */

src/include_hd/oclh_hd_squad.clh

@@ -0,0 +1,145 @@
+/*
+ * oclh_hd_squad.clh
+ *      Author: havock
+ */
+#ifndef _OCLH_OCL_COMPILER_
+#define __global
+#define __private
+#define __kernel
+#endif /* _OCLH_OCL_COMPILER_ */
+
+#ifndef OPENCL_INC_OCLH_SQUAD_COMMON_HEADER_CLH_DECLS_
+#define OPENCL_INC_OCLH_SQUAD_COMMON_HEADER_CLH_DECLS_ 1
+
+#if defined(_OCLH_OCL_HOST_HEADERS_) || defined(_OCLH_OCL_COMPILER_)
+#pragma pack(push,1)
+typedef struct _SQT_SQUAD_VARIABLES {
+    int32_t  i32Cmd; /* command */
+    uint32_t u32Cnt; /* counter of begun tasks, used as task index */
+    __global volatile uint32_t* pu32RptVals;
+} _SQT_SQUAD_VARS;
+#pragma pack(pop)
+#endif /* defined(_OCLH_OCL_HOST_HEADERS_) || defined(_OCLH_OCL_COMPILER_) */
+
+#ifdef _OCLH_OCL_COMPILER_
+
+#define _SQCMD_DISMISS  (-1)
+#define _SQCMD_WAIT     0
+
+#pragma pack(push,1)
+typedef struct _SQT_REPORT_BIT_AND_DATA {
+    uint64_t u64ValIdx;
+    uint64_t u64NofVals;
+    uint32_t u32Bit;
+    uint32_t u32FullVal;
+    uint32_t u32LastFullVal;
+} _SQT_RPT_BIT_N_DAT;
+typedef struct _SQT_TROOPER_TECHNICAL_DATA {
+              uint64_t u64MSN;      /* Military Service Number,
+                                       the unique overall trooper index */
+              uint64_t u64Role;     /* trooper index in squad */
+              uint64_t u64SquadNum; /* squad index */
+              uint64_t u64RegSz;    /* number of one squad */
+              uint64_t u64TrBufSzB; /* size of a trooper memory in bytes */
+    __global  uint8_t* pu8TrBuf;    /* pointer to a trooper memory in bytes */
+    __global  uint8_t* pu8SqZeroBuf;/* pointer to memory
+                                       of the zeroth trooper in squad */
+    _SQT_RPT_BIT_N_DAT RptBitNDat;
+} _SQT_TROOPER_TECHDAT;
+#pragma pack(pop)
+
+_SQT_TROOPER_TECHDAT _sqf_declPersRecWith(
+    __private const uint64_t u64MSN,
+    __private const uint64_t u64RegSz,
+    __global  uint8_t* const d_pu8TrooperBuf,
+    __private const uint64_t u64TrBufSzB);
+int32_t _sqf_giveCmdNRstRpt(
+    __global _SQT_SQUAD_VARS* const pSqDat,
+    __private const _SQT_RPT_BIT_N_DAT RptBitNDat,
+    __private const int32_t i32Cmd);
+uint32_t _sqf_setPrivSeqBitToOneLE32(__private const uint64_t u64N);
+int32_t _sqf_getRptBit(
+    __global const _SQT_SQUAD_VARS* const pSqDat,
+    __private _SQT_RPT_BIT_N_DAT RptBitNDat);
+int32_t _sqf_setBitToOne(
+    __global volatile uint32_t* pu32RptVals,
+    __private const _SQT_RPT_BIT_N_DAT RptBitNDat);
+
+#endif /* _OCLH_OCL_COMPILER_ */
+#endif /* OPENCL_INC_OCLH_SQUAD_COMMON_HEADER_CLH_DECLS_ */
+
+#ifdef _OCLH_OCL_COMPILER_
+#ifndef OPENCL_INC_OCLH_SQUAD_COMMON_HEADER_CLH_IMPLS_
+#define OPENCL_INC_OCLH_SQUAD_COMMON_HEADER_CLH_IMPLS_ 1
+
+_SQT_TROOPER_TECHDAT _sqf_declPersRecWith(
+    __private const uint64_t u64MSN,
+    __private const uint64_t u64RegSz,
+    __global uint8_t* const d_pu8TrooperBuf,
+    __private const uint64_t u64TrBufSzB) {
+    _SQT_TROOPER_TECHDAT PersRec;
+    PersRec.u64MSN=u64MSN;
+    PersRec.u64RegSz=u64RegSz;
+    PersRec.u64Role=u64MSN%u64RegSz;
+    PersRec.u64SquadNum=u64MSN/u64RegSz;
+    PersRec.pu8SqZeroBuf=
+        &d_pu8TrooperBuf[PersRec.u64SquadNum*u64TrBufSzB*u64RegSz];
+    PersRec.pu8TrBuf=&d_pu8TrooperBuf[u64MSN*u64TrBufSzB];
+    PersRec.u64TrBufSzB	=u64TrBufSzB;
+    {
+        PersRec.RptBitNDat.u64NofVals=
+            (u64RegSz%32)?((u64RegSz/32)+1):(u64RegSz/32);
+        PersRec.RptBitNDat.u64ValIdx=PersRec.u64Role/32;
+        PersRec.RptBitNDat.u32Bit=
+            _sqf_setPrivSeqBitToOneLE32(PersRec.u64Role%32);
+        PersRec.RptBitNDat.u32FullVal=0;
+        PersRec.RptBitNDat.u32FullVal=(~(PersRec.RptBitNDat.u32FullVal));
+        PersRec.RptBitNDat.u32LastFullVal=0;
+        if(u64RegSz%32) {
+            uint64_t i=0;
+            for(i=0;i<(u64RegSz%32);i++)
+                PersRec.RptBitNDat.u32LastFullVal=
+                    PersRec.RptBitNDat.u32LastFullVal^
+                        _sqf_setPrivSeqBitToOneLE32(i);
+        } else PersRec.RptBitNDat.u32LastFullVal=PersRec.RptBitNDat.u32FullVal;
+    }
+    return(PersRec);
+}
+
+int32_t _sqf_giveCmdNRstRpt(
+    __global _SQT_SQUAD_VARS* const pSqDat,
+    __private const _SQT_RPT_BIT_N_DAT RptBitNDat,
+    __private const int32_t i32Cmd) {
+    __private uint64_t i=0;
+    atomic_xchg(&pSqDat->i32Cmd,i32Cmd);
+    atomic_xchg(&pSqDat->u32Cnt,0);
+    for(i=0;i<RptBitNDat.u64NofVals;i++)
+        atomic_xchg(&pSqDat->pu32RptVals[i],0);
+    return(0);
+}
+
+uint32_t _sqf_setPrivSeqBitToOneLE32(__private const uint64_t u64N) {
+    __private uint32_t u32NewS=0;
+    __private uint8_t* pu8B=&(((__private uint8_t*) &u32NewS)[u64N/8]);
+    *pu8B=1<<(u64N%8);
+    return(u32NewS);
+}
+
+int32_t _sqf_getRptBit(
+    __global const _SQT_SQUAD_VARS* const pSqDat,
+    __private _SQT_RPT_BIT_N_DAT	RptBitNDat) {
+    __private uint32_t u32Val;
+    u32Val=pSqDat->pu32RptVals[RptBitNDat.u64ValIdx];
+    if((u32Val & RptBitNDat.u32Bit)==0) return(0);
+    else return(1);
+}
+
+int32_t _sqf_setBitToOne(
+    __global volatile uint32_t* pu32RptVals,
+    __private const _SQT_RPT_BIT_N_DAT RptBitNDat) {
+    atomic_or(&pu32RptVals[RptBitNDat.u64ValIdx],RptBitNDat.u32Bit);
+    return(0);
+}
+
+#endif /* OPENCL_INC_OCLH_SQUAD_COMMON_HEADER_CLH_IMPLS_ */
+#endif /* _OCLH_OCL_COMPILER_ */

src/include_hd/oclh_hd_srr.clh

@@ -0,0 +1,64 @@
+/*
+ * oclh_hd_srr.clh
+ *      Author: havock
+ */
+#ifndef _OCLH_OCL_COMPILER_
+#define __global
+#define __private
+#define __kernel
+#endif /* _OCLH_OCL_COMPILER_ */
+
+#if defined(_OCLH_OCL_HOST_HEADERS_) || defined(_OCLH_OCL_COMPILER_)
+#ifndef OCLH_HD_SRR_CLH_DEFS_
+#define OCLH_HD_SRR_CLH_DEFS_ 1
+#include <oclh_hd_std_types.clh>
+
+int32_t  _ghdf_sumFromAtoB_i32(__private const int32_t a,
+                               __private const int32_t b);
+uint64_t _ghdf_sumFromAtoB_u64(__private const uint64_t a,
+                               __private const uint64_t b);
+int32_t  _ghdf_getMaxOfInt32s( __private const int32_t* const pi32Vals,
+                               __private const uint64_t       u64NofVals);
+int32_t  _ghdf_getMaxOf2Int32s(__private const int32_t a,
+                               __private const int32_t b);
+
+#endif /* OCLH_HD_SRR_CLH_DEFS_ */
+#endif /* defined(_OCLH_OCL_HOST_HEADERS_) || defined(_OCLH_OCL_COMPILER_) */
+
+
+
+#if defined(_OCLH_OCL_HOST_ALGORITHMS_) || defined(_OCLH_OCL_COMPILER_)
+#ifndef OCLH_HD_SRR_CLH_DECLS_
+#define OCLH_HD_SRR_CLH_DECLS_ 1
+
+int32_t _ghdf_sumFromAtoB_i32(__private const int32_t a,
+                              __private const int32_t b) {
+    const int32_t d=b-a, s=b+a;
+    return((d%2)? /* i assume that -1 is TRUE ))) */
+            (s*((d+1)/2)):
+            ((s-1)*(d/2)+b));
+}
+
+uint64_t _ghdf_sumFromAtoB_u64(__private const uint64_t a,
+                               __private const uint64_t b) {
+    const uint64_t d=b-a, s=b+a;
+    return((d%2ul)? /* i assume that -1 is TRUE ))) */
+            (s*((d+1ul)/2ul)):
+            ((s-1ul)*(d/2ul)+b));
+}
+
+int32_t _ghdf_getMaxOfInt32s(__private const int32_t* const pi32Vals,
+                             __private const uint64_t       u64NofVals) {
+    __private int32_t i32Max=pi32Vals[0];
+    __private uint64_t i=1ul;
+    for(i=1ul; i<u64NofVals; i++) if(i32Max<pi32Vals[i]) i32Max=pi32Vals[i];
+    return(i32Max);
+}
+
+int32_t _ghdf_getMaxOf2Int32s(__private const int32_t a,
+                              __private const int32_t b) {
+    return((a>b)?a:b);
+}
+
+#endif /* OCLH_HD_SRR_CLH_DECLS_ */
+#endif /* defined(_OCLH_OCL_HOST_ALGORITHMS_) || defined(_OCLH_OCL_COMPILER_) */

src/include_hd/oclh_hd_std_types.clh

@@ -0,0 +1,30 @@
+/*
+ * oclh_hd_std_types.clh
+ *      Author: havock
+ */
+#ifndef OCLH_HD_STD_TYPES_CLH_
+#define OCLH_HD_STD_TYPES_CLH_ 1
+
+#ifndef _STDINT_H
+#define _STDINT_H 1
+typedef char                int8_t;
+typedef unsigned char       uint8_t;
+typedef short int           int16_t;
+typedef unsigned short int  uint16_t;
+typedef int                 int32_t;
+typedef unsigned int        uint32_t;
+typedef long int            int64_t;
+typedef unsigned long int   uint64_t;
+#endif /* _STDINT_H */
+
+#ifndef _STDFLT_H
+#define _STDFLT_H 1
+#ifdef _OCLH_OCL_COMPILER_
+typedef half                flt16_t;
+#endif /* _OCLH_OCL_COMPILER_ */
+typedef float               flt32_t;
+typedef double              flt64_t;
+typedef long double         flt128_t;
+#endif /* _STDFLT_H */
+
+#endif /* OCLH_HD_STD_TYPES_CLH_ */

src/oclh_cc.c

@@ -0,0 +1,114 @@
+/*
+ * oclh_cc.c
+ *
+ *  Created on: Jun 1, 2015
+ *      Author: havock
+ */
+#include <string.h>
+#include <oclh_cc.h>
+#include <oclh.h>
+
+int32_t main(int32_t argc, char *argv[]) {
+    int32_t err=0;
+    char *pcOutputPrefix  = NULL,
+         *pcOCLBuildOpts  = NULL,
+        **ppcSrcFilenames = NULL;
+    _GHT_LOG Log=_ghf_declLog();
+    _GHT_AWSS AWSs=_ghf_declAWSs(_GHE_LOG_ALL,
+                                 _GHE_BUILD_LOG_IN_SEPARATED_FILES);
+    {
+        uint64_t i=1ul;
+        for(i=1ul; i<(uint64_t)argc; i++) {
+            if(argv[i][0]=='-') { /* key-value */
+                if(argv[i][1]=='o') { /* output prefix */
+                    if(!argv[i][2]) { pcOutputPrefix=argv[++i]; }
+                    else { pcOutputPrefix=argv[i]+2ul; }
+                } else { /* compiler and linker options */
+                    err=__ghf_concatHeapStrAndCharPtr(&pcOCLBuildOpts,
+                                                       argv[i]);
+#if defined(__OCLH_BUILDER_FLAG) || defined(__OCLH_COMPILER_ONLY_FLAG)
+                    if((argv[i][1]=='D' || argv[i][1]=='I') && !argv[i][2])
+                        err|=__ghf_concatHeapStrAndCharPtr(&pcOCLBuildOpts,
+                                                           argv[++i]);
+                    if(!strncmp(argv[i],"-cl-std=",8ul) && !argv[i][8])
+                        err|=__ghf_concatHeapStrAndCharPtr(&pcOCLBuildOpts,
+                                                           argv[++i]);
+#endif /* defined(__OCLH_BUILDER_FLAG) || defined(__OCLH_COMPILER_ONLY_FLAG) */
+                }
+            } else err=_ghf_addCharPtrToCharPtrList(&ppcSrcFilenames,argv[i]);
+            if(err) {
+                fprintf(stderr,
+                        "Unable to add item \"%s\" to compiler options list\n",
+                        argv[i]);
+                _ghf_wipeCharPtrList(&ppcSrcFilenames);
+                _ghf_freeHostZ(&pcOCLBuildOpts);
+                return(err);
+            }
+        }
+    }
+    if(pcOutputPrefix) {
+        __ghf_removePreNPostSpacesFromCharPtr(pcOutputPrefix);
+        __ghf_replaceSpacesWithUnderscoreInCharPtr(pcOutputPrefix);
+    }
+    if(ppcSrcFilenames) {
+        uint64_t i=0ul;
+        for(i=0ul; ppcSrcFilenames[i]; i++)
+            __ghf_removePreNPostSpacesFromCharPtr(ppcSrcFilenames[i]);
+    } else {
+        _ghf_wipeCharPtrList(&ppcSrcFilenames);
+        _ghf_freeHostZ(&pcOCLBuildOpts);
+        fprintf(stderr, "No source files specified\n");
+        return(1);
+    }
+    err=_ghf_genrLog(&Log,_GHM_OCLH_BUILDER_LOG_FILENAME);
+    if(err) {
+        _ghf_wipeCharPtrList(&ppcSrcFilenames);
+        _ghf_freeHostZ(&pcOCLBuildOpts);
+        _ghf_wipeLog(&Log);
+        fprintf(stderr,"Unable to open the log file %s\n",
+                _GHM_OCLH_BUILDER_LOG_FILENAME);
+        return(err);
+    }
+    {
+        char** ppcSources=NULL;
+        uint64_t i=0ul;
+        while(ppcSrcFilenames[i]) {
+            err|=_ghf_addFileToCharPtrList(&ppcSources,ppcSrcFilenames[i],
+#if defined(__OCLH_BUILDER_FLAG) || defined(__OCLH_COMPILER_ONLY_FLAG)
+                                           8);
+#else
+                                           1);
+#endif /* defined(__OCLH_BUILDER_FLAG) || defined(__OCLH_COMPILER_ONLY_FLAG) */
+            i++;
+        }
+        err=_ghf_genrAWSs(&AWSs,Log,ppcSources,
+#if defined(__OCLH_BUILDER_FLAG) || defined(__OCLH_COMPILER_ONLY_FLAG)
+                          _GHE_HL_LISTINGS,
+#endif /* defined(__OCLH_BUILDER_FLAG) || defined(__OCLH_COMPILER_ONLY_FLAG) */
+#if defined(__OCLH_LINKER_ONLY_FLAG)
+                          _GHE_SEPARATED_OBJECTS,
+#endif /* defined(__OCLH_LINKER_ONLY_FLAG) */
+                          pcOCLBuildOpts
+#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) */
+                                          );
+        _ghf_wipeCharPtrList(&ppcSrcFilenames);
+        _ghf_wipeCharPtrList(&ppcSources);
+        _ghf_freeHostZ(&pcOCLBuildOpts);
+        if(err) fprintf(stderr,"Check the log file\n");
+    }
+    {
+        uint64_t i=0ul;
+        for(i=0ul; i<AWSs.u64NofWSs; i++) {
+            _ghf_saveWS_ProgramBinaries(AWSs.pWSet[i],pcOutputPrefix);
+        }
+    }
+    _ghf_wipeAWSs(&AWSs);
+    _ghf_wipeLog(&Log);
+    return(0);
+}

src/oclh_h_base_clapi_strings.c

@@ -0,0 +1,595 @@
+/*
+ * oclh_h_base_clapi_strings.c
+ *      Author: havock
+ */
+#include <oclh_h_base_clapi_strings.h>
+char* _ghf_CLAPIErrString(const cl_int clErr) {
+/* OpenCL 1.2 list */
+/*
+ * TODO: check OpenCL 2.2 codes
+ */
+    switch(clErr) {
+#ifdef CL_SUCCESS
+  case CL_SUCCESS:                  return("CL_SUCCESS");
+#endif
+#ifdef CL_DEVICE_NOT_FOUND
+  case CL_DEVICE_NOT_FOUND:         return("CL_DEVICE_NOT_FOUND");
+#endif
+#ifdef CL_DEVICE_NOT_AVAILABLE
+  case CL_DEVICE_NOT_AVAILABLE:     return("CL_DEVICE_NOT_AVAILABLE");
+#endif
+#ifdef CL_COMPILER_NOT_AVAILABLE
+  case CL_COMPILER_NOT_AVAILABLE:   return("CL_COMPILER_NOT_AVAILABLE");
+#endif
+#ifdef CL_MEM_OBJECT_ALLOCATION_FAILURE
+  case CL_MEM_OBJECT_ALLOCATION_FAILURE:
+                                    return("CL_MEM_OBJECT_ALLOCATION_FAILURE");
+#endif
+#ifdef CL_OUT_OF_RESOURCES
+  case CL_OUT_OF_RESOURCES:         return("CL_OUT_OF_RESOURCES");
+#endif
+#ifdef CL_OUT_OF_HOST_MEMORY
+  case CL_OUT_OF_HOST_MEMORY:       return("CL_OUT_OF_HOST_MEMORY");
+#endif
+#ifdef CL_PROFILING_INFO_NOT_AVAILABLE
+  case CL_PROFILING_INFO_NOT_AVAILABLE:
+                                    return("CL_PROFILING_INFO_NOT_AVAILABLE");
+#endif
+#ifdef CL_MEM_COPY_OVERLAP
+  case CL_MEM_COPY_OVERLAP:         return("CL_MEM_COPY_OVERLAP");
+#endif
+#ifdef CL_IMAGE_FORMAT_MISMATCH
+  case CL_IMAGE_FORMAT_MISMATCH:    return("CL_IMAGE_FORMAT_MISMATCH");
+#endif
+#ifdef CL_IMAGE_FORMAT_NOT_SUPPORTED
+  case CL_IMAGE_FORMAT_NOT_SUPPORTED:
+                                    return("CL_IMAGE_FORMAT_NOT_SUPPORTED");
+#endif
+#ifdef CL_BUILD_PROGRAM_FAILURE
+  case CL_BUILD_PROGRAM_FAILURE:    return("CL_BUILD_PROGRAM_FAILURE");
+#endif
+#ifdef CL_MAP_FAILURE
+  case CL_MAP_FAILURE:              return("CL_MAP_FAILURE");
+#endif
+#ifdef CL_MISALIGNED_SUB_BUFFER_OFFSET
+  case CL_MISALIGNED_SUB_BUFFER_OFFSET:
+                                    return("CL_MISALIGNED_SUB_BUFFER_OFFSET");
+#endif
+#ifdef CL_EXEC_STATUS_ERROR_FOR_EVENTS_IN_WAIT_LIST
+  case CL_EXEC_STATUS_ERROR_FOR_EVENTS_IN_WAIT_LIST:
+                         return("CL_EXEC_STATUS_ERROR_FOR_EVENTS_IN_WAIT_LIST");
+#endif
+#ifdef CL_COMPILE_PROGRAM_FAILURE
+  case CL_COMPILE_PROGRAM_FAILURE:  return("CL_COMPILE_PROGRAM_FAILURE");
+#endif
+#ifdef CL_LINKER_NOT_AVAILABLE
+  case CL_LINKER_NOT_AVAILABLE:     return("CL_LINKER_NOT_AVAILABLE");
+#endif
+#ifdef CL_LINK_PROGRAM_FAILURE
+  case CL_LINK_PROGRAM_FAILURE:     return("CL_LINK_PROGRAM_FAILURE");
+#endif
+#ifdef CL_DEVICE_PARTITION_FAILED
+  case CL_DEVICE_PARTITION_FAILED:  return("CL_DEVICE_PARTITION_FAILED");
+#endif
+#ifdef CL_KERNEL_ARG_INFO_NOT_AVAILABLE
+  case CL_KERNEL_ARG_INFO_NOT_AVAILABLE:
+                                    return("CL_KERNEL_ARG_INFO_NOT_AVAILABLE");
+#endif
+#ifdef CL_INVALID_VALUE
+  case CL_INVALID_VALUE:            return("CL_INVALID_VALUE");
+#endif
+#ifdef CL_INVALID_DEVICE_TYPE
+  case CL_INVALID_DEVICE_TYPE:      return("CL_INVALID_DEVICE_TYPE");
+#endif
+#ifdef CL_INVALID_PLATFORM
+  case CL_INVALID_PLATFORM:         return("CL_INVALID_PLATFORM");
+#endif
+#ifdef CL_INVALID_DEVICE
+  case CL_INVALID_DEVICE:           return("CL_INVALID_DEVICE");
+#endif
+#ifdef CL_INVALID_CONTEXT
+  case CL_INVALID_CONTEXT:          return("CL_INVALID_CONTEXT");
+#endif
+#ifdef CL_INVALID_QUEUE_PROPERTIES
+  case CL_INVALID_QUEUE_PROPERTIES: return("CL_INVALID_QUEUE_PROPERTIES");
+#endif
+#ifdef CL_INVALID_COMMAND_QUEUE
+  case CL_INVALID_COMMAND_QUEUE:    return("CL_INVALID_COMMAND_QUEUE");
+#endif
+#ifdef CL_INVALID_HOST_PTR
+  case CL_INVALID_HOST_PTR:         return("CL_INVALID_HOST_PTR");
+#endif
+#ifdef CL_INVALID_MEM_OBJECT
+  case CL_INVALID_MEM_OBJECT:       return("CL_INVALID_MEM_OBJECT");
+#endif
+#ifdef CL_INVALID_IMAGE_FORMAT_DESCRIPTOR
+  case CL_INVALID_IMAGE_FORMAT_DESCRIPTOR:
+                                   return("CL_INVALID_IMAGE_FORMAT_DESCRIPTOR");
+#endif
+#ifdef CL_INVALID_IMAGE_SIZE
+  case CL_INVALID_IMAGE_SIZE:       return("CL_INVALID_IMAGE_SIZE");
+#endif
+#ifdef CL_INVALID_SAMPLER
+  case CL_INVALID_SAMPLER:          return("CL_INVALID_SAMPLER");
+#endif
+#ifdef CL_INVALID_BINARY
+  case CL_INVALID_BINARY:           return("CL_INVALID_BINARY");
+#endif
+#ifdef CL_INVALID_BUILD_OPTIONS
+  case CL_INVALID_BUILD_OPTIONS:    return("CL_INVALID_BUILD_OPTIONS");
+#endif
+#ifdef CL_INVALID_PROGRAM
+  case CL_INVALID_PROGRAM:          return("CL_INVALID_PROGRAM");
+#endif
+#ifdef CL_INVALID_PROGRAM_EXECUTABLE
+  case CL_INVALID_PROGRAM_EXECUTABLE:
+                                    return("CL_INVALID_PROGRAM_EXECUTABLE");
+#endif
+#ifdef CL_INVALID_KERNEL_NAME
+  case CL_INVALID_KERNEL_NAME:      return("CL_INVALID_KERNEL_NAME");
+#endif
+#ifdef CL_INVALID_KERNEL_DEFINITION
+  case CL_INVALID_KERNEL_DEFINITION:return("CL_INVALID_KERNEL_DEFINITION");
+#endif
+#ifdef CL_INVALID_KERNEL
+  case CL_INVALID_KERNEL:           return("CL_INVALID_KERNEL");
+#endif
+#ifdef CL_INVALID_ARG_INDEX
+  case CL_INVALID_ARG_INDEX:        return("CL_INVALID_ARG_INDEX");
+#endif
+#ifdef CL_INVALID_ARG_VALUE
+  case CL_INVALID_ARG_VALUE:        return("CL_INVALID_ARG_VALUE");
+#endif
+#ifdef CL_INVALID_ARG_SIZE
+  case CL_INVALID_ARG_SIZE:         return("CL_INVALID_ARG_SIZE");
+#endif
+#ifdef CL_INVALID_KERNEL_ARGS
+  case CL_INVALID_KERNEL_ARGS:      return("CL_INVALID_KERNEL_ARGS");
+#endif
+#ifdef CL_INVALID_WORK_DIMENSION
+  case CL_INVALID_WORK_DIMENSION:   return("CL_INVALID_WORK_DIMENSION");
+#endif
+#ifdef CL_INVALID_WORK_GROUP_SIZE
+  case CL_INVALID_WORK_GROUP_SIZE:  return("CL_INVALID_WORK_GROUP_SIZE");
+#endif
+#ifdef CL_INVALID_WORK_ITEM_SIZE
+  case CL_INVALID_WORK_ITEM_SIZE:   return("CL_INVALID_WORK_ITEM_SIZE");
+#endif
+#ifdef CL_INVALID_GLOBAL_OFFSET
+  case CL_INVALID_GLOBAL_OFFSET:    return("CL_INVALID_GLOBAL_OFFSET");
+#endif
+#ifdef CL_INVALID_EVENT_WAIT_LIST
+  case CL_INVALID_EVENT_WAIT_LIST:  return("CL_INVALID_EVENT_WAIT_LIST");
+#endif
+#ifdef CL_INVALID_EVENT
+  case CL_INVALID_EVENT:            return("CL_INVALID_EVENT");
+#endif
+#ifdef CL_INVALID_OPERATION
+  case CL_INVALID_OPERATION:        return("CL_INVALID_OPERATION");
+#endif
+#ifdef CL_INVALID_GL_OBJECT
+  case CL_INVALID_GL_OBJECT:        return("CL_INVALID_GL_OBJECT");
+#endif
+#ifdef CL_INVALID_BUFFER_SIZE
+  case CL_INVALID_BUFFER_SIZE:      return("CL_INVALID_BUFFER_SIZE");
+#endif
+#ifdef CL_INVALID_MIP_LEVEL
+  case CL_INVALID_MIP_LEVEL:        return("CL_INVALID_MIP_LEVEL");
+#endif
+#ifdef CL_INVALID_GLOBAL_WORK_SIZE
+  case CL_INVALID_GLOBAL_WORK_SIZE: return("CL_INVALID_GLOBAL_WORK_SIZE");
+#endif
+#ifdef CL_INVALID_PROPERTY
+  case CL_INVALID_PROPERTY:         return("CL_INVALID_PROPERTY");
+#endif
+#ifdef CL_INVALID_IMAGE_DESCRIPTOR
+  case CL_INVALID_IMAGE_DESCRIPTOR: return("CL_INVALID_IMAGE_DESCRIPTOR");
+#endif
+#ifdef CL_INVALID_COMPILER_OPTIONS
+  case CL_INVALID_COMPILER_OPTIONS: return("CL_INVALID_COMPILER_OPTIONS");
+#endif
+#ifdef CL_INVALID_LINKER_OPTIONS
+  case CL_INVALID_LINKER_OPTIONS:   return("CL_INVALID_LINKER_OPTIONS");
+#endif
+#ifdef CL_INVALID_DEVICE_PARTITION_COUNT
+  case CL_INVALID_DEVICE_PARTITION_COUNT:
+                                    return("CL_INVALID_DEVICE_PARTITION_COUNT");
+#endif
+  case -9999:                       return("NVIDIA_CL_SEGMENTATION_FAULT");
+  default:                          return("Unknown OpenCL error");
+    }
+}
+
+char* _ghf_PlatfPropString(const cl_uint clProp) {
+/* OpenCL 2.2 list */
+    switch(clProp) {
+#ifdef CL_PLATFORM_PROFILE
+  case CL_PLATFORM_PROFILE:    return("Profile");
+#endif
+#ifdef CL_PLATFORM_VERSION
+  case CL_PLATFORM_VERSION:    return("Version");
+#endif
+#ifdef CL_PLATFORM_NAME
+  case CL_PLATFORM_NAME:       return("Name");
+#endif
+#ifdef CL_PLATFORM_VENDOR
+  case CL_PLATFORM_VENDOR:     return("Vendor");
+#endif
+#ifdef CL_PLATFORM_EXTENSIONS
+  case CL_PLATFORM_EXTENSIONS: return("Extensions");
+#endif
+#ifdef CL_PLATFORM_HOST_TIMER_RESOLUTION
+  case CL_PLATFORM_HOST_TIMER_RESOLUTION:
+      return("The resolution of the host timer");
+#endif
+        default: return("Unknown OpenCL platform property!");
+    }
+}
+
+char* _ghf_DevPropSuffixString(const cl_device_info clDevInf,
+                               const _GHE_RECTYPE   rt) {
+/* OpenCL 2.2 list */
+    switch(clDevInf) {
+#ifdef CL_DEVICE_TYPE
+        case CL_DEVICE_TYPE: return(rt?"Type:":"");
+#endif
+#ifdef CL_DEVICE_VENDOR_ID
+        case CL_DEVICE_VENDOR_ID: return(rt?"Vendor ID:":"");
+#endif
+#ifdef CL_DEVICE_MAX_COMPUTE_UNITS
+  case CL_DEVICE_MAX_COMPUTE_UNITS:
+      return(rt?"Max compute units (work-groups):":"");
+#endif
+#ifdef CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS
+  case CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS:
+      return(rt?"Max work-item dimensions:":"");
+#endif
+#ifdef CL_DEVICE_MAX_WORK_ITEM_SIZES
+  case CL_DEVICE_MAX_WORK_ITEM_SIZES: return(rt?"Max work-item sizes:":"");
+#endif
+#ifdef CL_DEVICE_MAX_WORK_GROUP_SIZE
+  case CL_DEVICE_MAX_WORK_GROUP_SIZE:
+      return(rt?"Max work-group size:":"work-items");
+#endif
+#ifdef CL_DEVICE_PREFERRED_VECTOR_WIDTH_CHAR
+  case CL_DEVICE_PREFERRED_VECTOR_WIDTH_CHAR:
+      return(rt?" Preferred vector size | char  :":"");
+#endif
+#ifdef CL_DEVICE_PREFERRED_VECTOR_WIDTH_SHORT
+  case CL_DEVICE_PREFERRED_VECTOR_WIDTH_SHORT:
+      return(rt?" Preferred vector size | short :":"");
+#endif
+#ifdef CL_DEVICE_PREFERRED_VECTOR_WIDTH_INT
+  case CL_DEVICE_PREFERRED_VECTOR_WIDTH_INT:
+      return(rt?" Preferred vector size | int   :":"");
+#endif
+#ifdef CL_DEVICE_PREFERRED_VECTOR_WIDTH_LONG
+  case CL_DEVICE_PREFERRED_VECTOR_WIDTH_LONG:
+      return(rt?" Preferred vector size | long  :":"");
+#endif
+#ifdef CL_DEVICE_PREFERRED_VECTOR_WIDTH_FLOAT
+  case CL_DEVICE_PREFERRED_VECTOR_WIDTH_FLOAT:
+      return(rt?" Preferred vector size | float :":"");
+#endif
+#ifdef CL_DEVICE_PREFERRED_VECTOR_WIDTH_DOUBLE
+  case CL_DEVICE_PREFERRED_VECTOR_WIDTH_DOUBLE:
+      return(rt?" Preferred vector size | double:":"");
+#endif
+#ifdef CL_DEVICE_PREFERRED_VECTOR_WIDTH_HALF
+  case CL_DEVICE_PREFERRED_VECTOR_WIDTH_HALF:
+      return(rt?" Preferred vector size | half  :":"");
+#endif
+#ifdef CL_DEVICE_NATIVE_VECTOR_WIDTH_CHAR
+  case CL_DEVICE_NATIVE_VECTOR_WIDTH_CHAR:
+      return(rt?" Native vector size | char  :":"");
+#endif
+#ifdef CL_DEVICE_NATIVE_VECTOR_WIDTH_SHORT
+  case CL_DEVICE_NATIVE_VECTOR_WIDTH_SHORT:
+      return(rt?" Native vector size | short :":"");
+#endif
+#ifdef CL_DEVICE_NATIVE_VECTOR_WIDTH_INT
+  case CL_DEVICE_NATIVE_VECTOR_WIDTH_INT:
+      return(rt?" Native vector size | int   :":"");
+#endif
+#ifdef CL_DEVICE_NATIVE_VECTOR_WIDTH_LONG
+  case CL_DEVICE_NATIVE_VECTOR_WIDTH_LONG:
+      return(rt?" Native vector size | long  :":"");
+#endif
+#ifdef CL_DEVICE_NATIVE_VECTOR_WIDTH_FLOAT
+  case CL_DEVICE_NATIVE_VECTOR_WIDTH_FLOAT:
+      return(rt?" Native vector size | float :":"");
+#endif
+#ifdef CL_DEVICE_NATIVE_VECTOR_WIDTH_DOUBLE
+  case CL_DEVICE_NATIVE_VECTOR_WIDTH_DOUBLE:
+      return(rt?" Native vector size | double:":"");
+#endif
+#ifdef CL_DEVICE_NATIVE_VECTOR_WIDTH_HALF
+  case CL_DEVICE_NATIVE_VECTOR_WIDTH_HALF:
+      return(rt?" Native vector size | half  :":"");
+#endif
+#ifdef CL_DEVICE_MAX_CLOCK_FREQUENCY
+  case CL_DEVICE_MAX_CLOCK_FREQUENCY: return(rt?"Max clock frequency:":"MHz");
+#endif
+#ifdef CL_DEVICE_ADDRESS_BITS
+  case CL_DEVICE_ADDRESS_BITS:        return(rt?"Address space size:":"bits");
+#endif
+#ifdef CL_DEVICE_MAX_MEM_ALLOC_SIZE
+  case CL_DEVICE_MAX_MEM_ALLOC_SIZE:return(rt?"Max memory allocation:":"bytes");
+#endif
+#ifdef CL_DEVICE_IMAGE_SUPPORT
+  case CL_DEVICE_IMAGE_SUPPORT:     return(rt?"Image support:":"");
+#endif
+#ifdef CL_DEVICE_MAX_READ_IMAGE_ARGS
+  case CL_DEVICE_MAX_READ_IMAGE_ARGS:
+      return(rt?" Image support | "
+                "Max number of image objects arguments for reading:":"");
+#endif
+#ifdef CL_DEVICE_MAX_WRITE_IMAGE_ARGS
+  case CL_DEVICE_MAX_WRITE_IMAGE_ARGS:
+      return(rt?" Image support | "
+                "Max number of image objects arguments for writing:":"");
+#endif
+#ifdef CL_DEVICE_MAX_READ_WRITE_IMAGE_ARGS
+  case CL_DEVICE_MAX_READ_WRITE_IMAGE_ARGS:
+      return(rt?" Image support | "
+                "Max number of image objects arguments for reading and writing:"
+                :"");
+#endif
+#ifdef CL_DEVICE_IMAGE2D_MAX_WIDTH
+  case CL_DEVICE_IMAGE2D_MAX_WIDTH:
+      return(rt?" Image support | Max 2D-image width:":"pixels");
+#endif
+#ifdef CL_DEVICE_IMAGE2D_MAX_HEIGHT
+  case CL_DEVICE_IMAGE2D_MAX_HEIGHT:
+      return(rt?" Image support | Max 2D-image height:":"pixels");
+#endif
+#ifdef CL_DEVICE_IMAGE3D_MAX_WIDTH
+  case CL_DEVICE_IMAGE3D_MAX_WIDTH:
+      return(rt?" Image support | Max 3D-image width:":"pixels");
+#endif
+#ifdef CL_DEVICE_IMAGE3D_MAX_HEIGHT
+  case CL_DEVICE_IMAGE3D_MAX_HEIGHT:
+      return(rt?" Image support | Max 3D-image height:":"pixels");
+#endif
+#ifdef CL_DEVICE_IMAGE3D_MAX_DEPTH
+  case CL_DEVICE_IMAGE3D_MAX_DEPTH:
+      return(rt?" Image support | Max 3D-image depth:":"pixels");
+#endif
+#ifdef CL_DEVICE_IMAGE_MAX_BUFFER_SIZE
+  case CL_DEVICE_IMAGE_MAX_BUFFER_SIZE:
+      return(rt?" Image support | "
+                "Max number of pixels for 1D-image from buffer:":"pixels");
+#endif
+#ifdef CL_DEVICE_IMAGE_MAX_ARRAY_SIZE
+  case CL_DEVICE_IMAGE_MAX_ARRAY_SIZE:
+      return(rt?" Image support | "
+                "Max number of images in 1D or 2D image array:":"images");
+#endif
+#ifdef CL_DEVICE_MAX_SAMPLERS
+  case CL_DEVICE_MAX_SAMPLERS:
+      return(rt?" Image support | "
+                "Max number of samplers used in a kernel:":"samplers");
+#endif
+#ifdef CL_DEVICE_IMAGE_PITCH_ALIGNMENT
+  case CL_DEVICE_IMAGE_PITCH_ALIGNMENT:
+      return(rt?" Image support | "
+                "The row pitch alignment size for 2D images from a buffer:"
+                :"pixels");
+#endif
+#ifdef CL_DEVICE_IMAGE_BASE_ADDRESS_ALIGNMENT
+  case CL_DEVICE_IMAGE_BASE_ADDRESS_ALIGNMENT:
+      return(rt?" Image support | Min alignment of the host_ptr:":"pixels");
+#endif
+#ifdef CL_DEVICE_MAX_PIPE_ARGS
+  case CL_DEVICE_MAX_PIPE_ARGS:
+      return(rt?"Max number of pipes passed as arguments to a kernel:":"");
+#endif
+#ifdef CL_DEVICE_PIPE_MAX_ACTIVE_RESERVATIONS
+  case CL_DEVICE_PIPE_MAX_ACTIVE_RESERVATIONS:
+      return(rt?"Max number of active reservations for a pipe "
+                "per work-item in a kernel:":"");
+#endif
+#ifdef CL_DEVICE_PIPE_MAX_PACKET_SIZE
+  case CL_DEVICE_PIPE_MAX_PACKET_SIZE: return(rt?"Max size of pipe packet:"
+                                                 :"bytes");
+#endif
+#ifdef CL_DEVICE_MAX_PARAMETER_SIZE
+  case CL_DEVICE_MAX_PARAMETER_SIZE:
+      return(rt?"Max size of all arguments passed to a kernel:":"bytes");
+#endif
+#ifdef CL_DEVICE_MEM_BASE_ADDR_ALIGN
+  case CL_DEVICE_MEM_BASE_ADDR_ALIGN:
+      return(rt?"Alignment for sub-buffer offsets:":"bits");
+#endif
+#ifdef CL_DEVICE_SINGLE_FP_CONFIG
+  case CL_DEVICE_SINGLE_FP_CONFIG:
+      return(rt?"Supported single precision floating-point capabilities:":"");
+#endif
+#ifdef CL_DEVICE_DOUBLE_FP_CONFIG
+  case CL_DEVICE_DOUBLE_FP_CONFIG:
+      return(rt?"Supported double precision floating-point capabilities:":"");
+#endif
+#ifdef CL_DEVICE_GLOBAL_MEM_CACHE_TYPE
+  case CL_DEVICE_GLOBAL_MEM_CACHE_TYPE:
+      return(rt?"Global memory cache support:":"");
+#endif
+#ifdef CL_DEVICE_GLOBAL_MEM_CACHELINE_SIZE
+  case CL_DEVICE_GLOBAL_MEM_CACHELINE_SIZE:
+      return(rt?"Size of global memory cache line:":"bytes");
+#endif
+#ifdef CL_DEVICE_GLOBAL_MEM_CACHE_SIZE
+  case CL_DEVICE_GLOBAL_MEM_CACHE_SIZE:
+      return(rt?"Size of global memory cache:":"bytes");
+#endif
+#ifdef CL_DEVICE_GLOBAL_MEM_SIZE
+  case CL_DEVICE_GLOBAL_MEM_SIZE:
+      return(rt?"Size of global device memory:":"bytes");
+#endif
+#ifdef CL_DEVICE_MAX_CONSTANT_BUFFER_SIZE
+  case CL_DEVICE_MAX_CONSTANT_BUFFER_SIZE:
+      return(rt?"Max size of a constant buffer allocation:":"bytes");
+#endif
+#ifdef CL_DEVICE_MAX_CONSTANT_ARGS
+  case CL_DEVICE_MAX_CONSTANT_ARGS:
+      return(rt?"Max number of arguments declared with "
+                "the __constant qualifier in a kernel:":"arguments");
+#endif
+#ifdef CL_DEVICE_MAX_GLOBAL_VARIABLE_SIZE
+  case CL_DEVICE_MAX_GLOBAL_VARIABLE_SIZE:
+      return(rt?"Max size of storage that may be allocated for any single "
+                "variable in program scope or inside a function in OpenCL C "
+                "declared in the global address space:":"bytes");
+#endif
+#ifdef CL_DEVICE_GLOBAL_VARIABLE_PREFERRED_TOTAL_SIZE
+  case CL_DEVICE_GLOBAL_VARIABLE_PREFERRED_TOTAL_SIZE:
+      return(rt?"Maximum preferred total size of all program variables "
+                "in the global address space:":"bytes");
+#endif
+#ifdef CL_DEVICE_LOCAL_MEM_TYPE
+  case CL_DEVICE_LOCAL_MEM_TYPE:return(rt?"Type of local memory supported:":"");
+#endif
+#ifdef CL_DEVICE_LOCAL_MEM_SIZE
+  case CL_DEVICE_LOCAL_MEM_SIZE: return(rt?"Size of local memory region:":"");
+#endif
+#ifdef CL_DEVICE_ERROR_CORRECTION_SUPPORT
+  case CL_DEVICE_ERROR_CORRECTION_SUPPORT:
+      return(rt?"Error correction support:":"");
+#endif
+#ifdef CL_DEVICE_PROFILING_TIMER_RESOLUTION
+  case CL_DEVICE_PROFILING_TIMER_RESOLUTION:
+      return(rt?"Resolution of device timer:":"nanoseconds");
+#endif
+#ifdef CL_DEVICE_ENDIAN_LITTLE
+  case CL_DEVICE_ENDIAN_LITTLE: return(rt?"Endianness:":"");
+#endif
+#ifdef CL_DEVICE_AVAILABLE
+  case CL_DEVICE_AVAILABLE: return(rt?"Is device available:":"");
+#endif
+#ifdef CL_DEVICE_COMPILER_AVAILABLE
+  case CL_DEVICE_COMPILER_AVAILABLE:
+      return(rt?"Is device compiler available:":"");
+#endif
+#ifdef CL_DEVICE_LINKER_AVAILABLE
+  case CL_DEVICE_LINKER_AVAILABLE: return(rt?"Is device linker available:":"");
+#endif
+  /*
+   * TODO: CL_DEVICE_IL_VERSION
+   */
+#ifdef CL_DEVICE_EXECUTION_CAPABILITIES
+  case CL_DEVICE_EXECUTION_CAPABILITIES:
+      return(rt?"Execution capabilities of the device:":"");
+#endif
+#ifdef CL_DEVICE_QUEUE_PROPERTIES /* 2.2 abs */
+  case CL_DEVICE_QUEUE_PROPERTIES:
+      return(rt?"Command-queue properties supported by the device:":"");
+#endif
+#ifdef CL_DEVICE_QUEUE_ON_HOST_PROPERTIES
+  case CL_DEVICE_QUEUE_ON_HOST_PROPERTIES:
+      return(rt?"Host command-queue properties supported by the device:":"");
+#endif
+#ifdef CL_DEVICE_QUEUE_ON_DEVICE_PROPERTIES
+  case CL_DEVICE_QUEUE_ON_DEVICE_PROPERTIES:
+      return(rt?"Device command-queue properties supported by the device":"");
+#endif
+#ifdef CL_DEVICE_QUEUE_ON_DEVICE_PREFERRED_SIZE
+  case CL_DEVICE_QUEUE_ON_DEVICE_PREFERRED_SIZE:
+      return(rt?"The preferred size of the device queue:":"bytes");
+#endif
+#ifdef CL_DEVICE_QUEUE_ON_DEVICE_MAX_SIZE
+  case CL_DEVICE_QUEUE_ON_DEVICE_MAX_SIZE:
+      return(rt?"The max size of the device queue:":"bytes");
+#endif
+#ifdef CL_DEVICE_MAX_ON_DEVICE_QUEUES
+  case CL_DEVICE_MAX_ON_DEVICE_QUEUES:
+      return(rt?"The maximum number of device queues per context:":"queues");
+#endif
+#ifdef CL_DEVICE_MAX_ON_DEVICE_EVENTS
+  case CL_DEVICE_MAX_ON_DEVICE_EVENTS:
+      return(rt?"The maximum number of events in use by a device queue:"
+                :"events");
+#endif
+#ifdef CL_DEVICE_BUILT_IN_KERNELS
+  case CL_DEVICE_BUILT_IN_KERNELS: return(rt?"Built-in kernels:":"");
+#endif
+#ifdef CL_DEVICE_PLATFORM
+  case CL_DEVICE_PLATFORM: return(rt?"Platform:":"");
+#endif
+#ifdef CL_DEVICE_NAME
+  case CL_DEVICE_NAME: return(rt?"Device name:":"");
+#endif
+#ifdef CL_DEVICE_VENDOR
+  case CL_DEVICE_VENDOR: return(rt?"Device vendor:":"");
+#endif
+#ifdef CL_DRIVER_VERSION
+  case CL_DRIVER_VERSION: return(rt?"Driver version:":"");
+#endif
+#ifdef CL_DEVICE_PROFILE
+  case CL_DEVICE_PROFILE: return(rt?"Device profile:":"");
+#endif
+#ifdef CL_DEVICE_VERSION
+  case CL_DEVICE_VERSION: return(rt?"Device version:":"");
+#endif
+#ifdef CL_DEVICE_OPENCL_C_VERSION
+  case CL_DEVICE_OPENCL_C_VERSION: return(rt?"Device OpenCL C version:":"");
+#endif
+#ifdef CL_DEVICE_EXTENSIONS
+  case CL_DEVICE_EXTENSIONS: return(rt?"Extensions:":"");
+#endif
+#ifdef CL_DEVICE_PRINTF_BUFFER_SIZE
+  case CL_DEVICE_PRINTF_BUFFER_SIZE:
+      return(rt?"Maximum size of the internal printf() buffer:":"bytes");
+#endif
+#ifdef CL_DEVICE_PREFERRED_INTEROP_USER_SYNC
+  case CL_DEVICE_PREFERRED_INTEROP_USER_SYNC:
+      return(rt?"Is device preferred for the user's responsibility for sync:"
+                :"");
+#endif
+#ifdef CL_DEVICE_PARENT_DEVICE
+  case CL_DEVICE_PARENT_DEVICE: return(rt?"The parent device:":"");
+#endif
+#ifdef CL_DEVICE_PARTITION_MAX_SUB_DEVICES
+  case CL_DEVICE_PARTITION_MAX_SUB_DEVICES:
+      return(rt?"Max number of sub-devices:":"");
+#endif
+    /*
+     * TODO: CL_DEVICE_PARTITION_PROPERTIES
+     * TODO: CL_DEVICE_PARTITION_AFFINITY_DOMAIN
+     * TODO: CL_DEVICE_PARTITION_TYPE
+     */
+#ifdef CL_DEVICE_REFERENCE_COUNT
+  case CL_DEVICE_REFERENCE_COUNT: return(rt?"Device reference count:":"");
+#endif
+#ifdef CL_DEVICE_SVM_CAPABILITIES
+  case CL_DEVICE_SVM_CAPABILITIES:
+      return(rt?"Shared virtual memory allocation types the device supports:"
+                :"");
+#endif
+#ifdef CL_DEVICE_PREFERRED_PLATFORM_ATOMIC_ALIGNMENT
+  case CL_DEVICE_PREFERRED_PLATFORM_ATOMIC_ALIGNMENT:
+      return(rt?"Preferred alignment for atomic types:":"bytes");
+#endif
+#ifdef CL_DEVICE_PREFERRED_GLOBAL_ATOMIC_ALIGNMENT
+  case CL_DEVICE_PREFERRED_GLOBAL_ATOMIC_ALIGNMENT:
+      return(rt?"Preferred alignment for atomic types to global memory:"
+                :"bytes");
+#endif
+#ifdef CL_DEVICE_PREFERRED_LOCAL_ATOMIC_ALIGNMENT
+  case CL_DEVICE_PREFERRED_LOCAL_ATOMIC_ALIGNMENT:
+      return(rt?"Preferred alignment for atomic types to local memory:"
+                :"bytes");
+#endif
+#ifdef CL_DEVICE_MAX_NUM_SUB_GROUPS
+  case CL_DEVICE_MAX_NUM_SUB_GROUPS:
+      return(rt?"Maximum number of sub-groups in a work-group that a device is "
+                "capable of executing on a single compute unit, for any given "
+                "kernel-instance running on the device:":"sub-groups");
+#endif
+#ifdef CL_DEVICE_SUB_GROUP_INDEPENDENT_FORWARD_PROGRESS
+  case CL_DEVICE_SUB_GROUP_INDEPENDENT_FORWARD_PROGRESS:
+      return(rt?"Does the device support "
+                "independent forward progress of sub-groups:":"");
+#endif
+        default: return(rt?"Unknown OpenCL device property!"
+                           :"Unknown measure!");
+    }
+}

src/oclh_h_base_dev_clapi_wrappers.c

@@ -0,0 +1,69 @@
+/*
+ * oclh_h_base_dev_clapi_wrappers.c
+ *      Author: havock
+ */
+
+#include <oclh_h_base_dev_clapi_wrappers.h>
+#include <oclh_h_base_defs.h>
+
+cl_bool _ghf_getDevInf_bool(const cl_device_id   clDev,
+                            const cl_device_info InfoVal,
+                                  cl_int*  const pclErr) {
+    cl_bool clbRes=0u; size_t szRes=0ul;
+    *pclErr=clGetDeviceInfo(clDev, InfoVal, sizeof(cl_bool),&clbRes,&szRes);
+    if(*pclErr) clbRes=(cl_bool)_GHM_UNDEFUINTVAL;
+    return(clbRes);
+}
+
+char* _ghf_getDevInf_charptr(const cl_device_id   clDev,
+                             const cl_device_info InfoVal,
+                                   cl_int*  const pclErr) {
+    size_t szRes=0ul; char* pcInf=NULL;
+    *pclErr=clGetDeviceInfo(clDev, InfoVal, 0ul, NULL, &szRes);
+    if(!*pclErr) {
+        if(szRes) {
+            pcInf=malloc(szRes);
+            if(pcInf) {
+                *pclErr=clGetDeviceInfo(clDev, InfoVal, szRes, pcInf, NULL);
+                if(*pclErr) { free(pcInf); pcInf=NULL; }
+            }
+        } else { pcInf=malloc(1ul); if(pcInf) pcInf[0]='\0'; }
+    }
+    return(pcInf);
+}
+
+cl_uint _ghf_getDevInf_cluint(const cl_device_id   clDev,
+                              const cl_device_info InfoVal,
+                                    cl_int*  const pclErr) {
+    cl_uint cluRes=0u; size_t szRes=0ul;
+    *pclErr=clGetDeviceInfo(clDev, InfoVal, sizeof(cl_uint), &cluRes, &szRes);
+    if(*pclErr) cluRes=(cl_uint)_GHM_UNDEFUINTVAL;
+    return(cluRes);
+}
+
+cl_ulong _ghf_getDevInf_clulong(const cl_device_id   clDev,
+                                const cl_device_info InfoVal,
+                                      cl_int*  const pclErr) {
+    cl_ulong Res=0ul; size_t szRes=0ul;
+    *pclErr=clGetDeviceInfo(clDev, InfoVal, sizeof(cl_ulong), &Res, &szRes);
+    if(*pclErr) Res=_GHM_UNDEFUINTVAL;
+    return(Res);
+}
+
+size_t _ghf_getDevInf_size(const cl_device_id   clDev,
+                           const cl_device_info InfoVal,
+                                 cl_int*  const pclErr) {
+    size_t Res=0ul; size_t szRes=0ul;
+    *pclErr=clGetDeviceInfo(clDev, InfoVal, sizeof(size_t), &Res, &szRes);
+    if(*pclErr) Res=_GHM_UNDEFUINTVAL;
+    return(Res);
+}
+
+uintptr_t _ghf_getDevInf_uintptr(const cl_device_id   clDev,
+                                 const cl_device_info InfoVal,
+                                       cl_int*  const pclErr) {
+    uintptr_t ptrRes=(uintptr_t)NULL; size_t szRes=0ul;
+    *pclErr=clGetDeviceInfo(clDev, InfoVal, sizeof(uintptr_t), &ptrRes, &szRes);
+    if(*pclErr) ptrRes=(uintptr_t)_GHM_UNDEFPTR;
+    return(ptrRes);
+}

src/oclh_h_base_log.c

@@ -0,0 +1,49 @@
+/*
+ * oclh_h_base_log.с
+ *      Author: havock
+ */
+#include <stdlib.h>
+#include <oclh_h_base_log.h>
+
+_GHT_LOG _ghf_declLog(void) {
+    _GHT_LOG log={ .pfOut=NULL, .pMtx=NULL };
+    return(log);
+}
+
+int32_t _ghf_genrLog(_GHT_LOG* const pLog, const char* const pcFileName) {
+    FILE*           pFile=NULL;
+    pthread_mutex_t EmptyMtx=PTHREAD_MUTEX_INITIALIZER,
+                    *pMtx=NULL ;
+    _ghf_wipeLog(pLog);
+    if(pcFileName) {
+        pFile=fopen(pcFileName,"w");
+        if(!pFile) { perror(""); return(_GHM_OPEN_FILE_ERROR); }
+    } else pFile=stdout;
+    pMtx=malloc(sizeof(pthread_mutex_t));
+    if(!pMtx) { fclose(pFile); perror(""); return(_GHM_HOST_MEMALLOC_ERROR); }
+    *pMtx=EmptyMtx;
+    pLog->pfOut=pFile;
+    pLog->pMtx=pMtx;
+    return(_GHM_OK);
+}
+
+int32_t _ghf_isLogValid(const _GHT_LOG Log) {
+    return(Log.pMtx && Log.pfOut);
+}
+
+int32_t _ghf_wipeLog(_GHT_LOG* const pLog) {
+    if(pLog) {
+        if(pLog->pMtx) {
+            pthread_mutex_lock(pLog->pMtx);
+            if(pLog->pfOut && pLog->pfOut!=stdout) fclose(pLog->pfOut);
+            pLog->pfOut=NULL;
+            pthread_mutex_unlock(pLog->pMtx);
+        } else {
+            if(pLog->pfOut && pLog->pfOut!=stdout) fclose(pLog->pfOut);
+            pLog->pfOut=NULL;
+        }
+        if(pLog->pMtx) { free(pLog->pMtx); pLog->pMtx=NULL; }
+        pLog->pMtx=NULL;
+    }
+    return(_GHM_OK);
+}

src/oclh_h_clapi_callbacks.c

@@ -0,0 +1,20 @@
+/*
+ * oclh_clapi_callbacks.c
+ *      Author: havock
+ */
+#include <stdio.h>
+#include <oclh_h_clapi_callbacks.h>
+
+void CL_CALLBACK _ghf_CtxEvent(const char*  pcErrInfo,
+                               const void*  pvPrivInfo,
+                                     size_t szPrivInfo,
+                                     void*  pvUserData) {
+    printf("Ctx callback | ");
+    if(pcErrInfo) printf("err: %s", pcErrInfo);
+    printf("priv_inf_addr:0x%lx ", (uint64_t)pvPrivInfo);
+    printf("priv_inf_sz: %f MiB/%f KiB/%lu B ",
+           ((float)szPrivInfo)/1024e0f/1024e0f, ((float)szPrivInfo)/1024e0f,
+           (uint64_t)szPrivInfo);
+    printf("user_data_addr:0x%lx\n", (uint64_t)pvUserData);
+    return;
+}

src/oclh_h_externals.c

@@ -0,0 +1,178 @@
+/*
+ * oclh_h_externals.c
+ *      Author: havock
+ */
+#include <stdio.h>
+#include <stdlib.h>
+#include <stdint.h>
+#include <string.h>
+#include <oclh_hd_std_types.clh>
+#include <oclh_h_settings.h>
+#include <oclh_h_externals.h>
+#include <oclh_h_errors.h>
+
+void* _ghf_wdcAllocHostBuf(_GHT_WRKSET wSet, const size_t sz) {
+    void* res=malloc(sz);
+    if(!res) {
+        if(wSet.pvDat) wSet.pfnDatCleaner(wSet);
+        return(NULL);
+    }
+    return(res);
+}
+
+void _ghf_freeHostZ(void* const ppPtr) {
+/* It looks like insanity but it is necessary. The problem is if you copy mem
+ * from GPU to host with clEnqueueReadBuffer() you get unobvious mapping between
+ * hostPointer and gpuPointer as a result freeing hostPointer without explicit
+ * zeroing it and following clReleaseMemObject() on GpuPointer causes errors
+ * like a corrupted linked-list or segmentation fault */
+    void* pRealPtr=(*((void**) ppPtr));
+    if(pRealPtr) { free(pRealPtr); *((void**)ppPtr)=NULL; }
+    return;
+}
+
+cl_mem _ghf_wdcAllocDevBuf(_GHT_WRKSET  wSet,
+                     const cl_mem_flags cmFlags,
+                     const size_t       sz,
+                           void*  const pvHostPtr) {
+    size_t szAct=0ul;
+    cl_mem res=clCreateBuffer(_ghf_getWS_Ctx(wSet), cmFlags, sz,
+                              pvHostPtr, &wSet.APIErr);
+    char pcLogMsg[_GHM_MAXLEN_OF_LOGMSG];
+    snprintf(pcLogMsg, _GHM_MAXLEN_OF_LOGMSG, "%s/clCreateBuffer", __func__);
+    if(_ghf_wdcChkWS_APIErr(wSet, pcLogMsg, _GHM_FL)) return(NULL);
+    wSet.APIErr=clGetMemObjectInfo(res, CL_MEM_SIZE,
+                                   sizeof(size_t), &szAct, NULL);
+    snprintf(pcLogMsg, _GHM_MAXLEN_OF_LOGMSG,
+             "%s/clGetMemObjectInfo/CL_MEM_SIZE", __func__);
+    _ghf_wdcChkWS_APIErr(wSet, pcLogMsg, _GHM_FL);
+    if(wSet.APIErr) res=NULL;
+    else (*wSet.pszMemBytes)+=szAct;
+    return(res);
+}
+
+void _ghf_freeDevZ(_GHT_WRKSET wSet, cl_mem* const pCLMem) {
+/* the current wrapper have made because of clReleaseMemObject() in contrast to
+ * free() did not check the NULL pointer and what was more important it was
+ * comme il faut to set a freed pointer to NULL */
+    wSet.pvDat=NULL;
+    if(*pCLMem) {
+        size_t szAct=0ul;
+        char pcLogMsg[_GHM_MAXLEN_OF_LOGMSG];
+        snprintf(pcLogMsg, _GHM_MAXLEN_OF_LOGMSG,
+                 "%s/clGetMemObjectInfo/CL_MEM_SIZE", __func__);
+        wSet.APIErr=clGetMemObjectInfo(*pCLMem, CL_MEM_SIZE,
+                                       sizeof(size_t), &szAct,NULL);
+        _ghf_wdcChkWS_APIErr(wSet,pcLogMsg, _GHM_FL);
+        wSet.APIErr=clReleaseMemObject(*pCLMem);
+        snprintf(pcLogMsg, _GHM_MAXLEN_OF_LOGMSG,
+                 "%s/clReleaseMemObject", __func__);
+        if(!_ghf_wdcChkWS_APIErr(wSet, pcLogMsg, _GHM_FL)) {
+            *pCLMem=NULL; (*wSet.pszMemBytes)-=szAct;
+        }
+    }
+    return;
+}
+
+int32_t _ghf_addCharPtrToCharPtrList(char*** const pppcLst,
+                                     const char* const pcStr) {
+    if(pppcLst && pcStr) {
+        char** ppcTmp=NULL;
+        uint64_t i=0, u64StrLen=strlen(pcStr);
+        if(*pppcLst) while((*pppcLst)[i]) i++;
+        else {
+            *pppcLst=malloc(sizeof(char*));
+            if(!*pppcLst) return(_GHM_HOST_MEMALLOC_ERROR);
+            (*pppcLst)[0]=NULL;
+        }
+        ppcTmp=realloc(*pppcLst,sizeof(char*)*(i+2ul));
+        if(!ppcTmp) return(_GHM_HOST_MEMALLOC_ERROR);
+        *pppcLst=ppcTmp;
+        (*pppcLst)[i]=NULL;
+        (*pppcLst)[i]=malloc(sizeof(char)*(u64StrLen+1));
+        if(!(*pppcLst)[i]) return(_GHM_HOST_MEMALLOC_ERROR);
+        strncpy((*pppcLst)[i],pcStr,u64StrLen);
+        (*pppcLst)[i+1]=NULL;
+        return(_GHM_OK);
+    } else return(_GHM_NULL_POINTER_RECEIVED_ERROR);
+}
+
+int32_t _ghf_addFileToCharPtrList(      char*** const pppcLst,
+                                  const char*   const pcFileName,
+                                  const uint64_t      u64Align) {
+    if(pppcLst && pcFileName) {
+        FILE* pFile=NULL;
+        size_t szFileLen=0ul;
+        char* pcSource=NULL;
+        pFile=fopen(pcFileName, "rb");
+        if(!pFile) {
+            fprintf(stderr, "Unable to open the source file %s\n", pcFileName);
+            perror("");
+            return(_GHM_OPEN_FILE_ERROR);
+        }
+        if(fseek(pFile,0L,SEEK_END)) {
+            fprintf(stderr,"Unable to seek the source file %s\n",pcFileName);
+            perror("");
+            fclose(pFile);
+            return(_GHM_SEEK_FILE_ERROR);
+        }
+        szFileLen=(size_t)ftell(pFile);
+        if(fseek(pFile,0L,SEEK_SET)) {
+            fprintf(stderr,"Unable to seek the source file %s\n",pcFileName);
+            perror("");
+            fclose(pFile);
+            return(_GHM_SEEK_FILE_ERROR);
+        }
+        pcSource=malloc((szFileLen%u64Align)?szFileLen+
+                        (u64Align-szFileLen%u64Align):szFileLen); /* align */
+        if(!pcSource) {
+            fprintf(stderr,"Unable to allocate host memory\n");
+            perror("");
+            fclose(pFile);
+            return(_GHM_HOST_MEMALLOC_ERROR);
+        }
+        if(fread(pcSource,szFileLen,1,pFile)!=1) {
+            fprintf(stderr,"Unable to read the source file %s\n",pcFileName);
+            perror("");
+            _ghf_freeHostZ(&pcSource);
+            fclose(pFile);
+            return(_GHM_READ_FILE_ERROR);
+        }
+        fclose(pFile);
+        {
+            char** ppcTmp=NULL;
+            uint64_t i=0ul;
+            if(*pppcLst) while((*pppcLst)[i]) i++;
+            if(!*pppcLst) {
+                *pppcLst=malloc(sizeof(char*));
+                if(!pcSource) {
+                    fprintf(stderr,"Unable to allocate host memory\n");
+                    perror("");
+                    _ghf_freeHostZ(pcSource);
+                    return(_GHM_HOST_MEMALLOC_ERROR);
+                }
+                (*pppcLst)[0]=NULL;
+            }
+            ppcTmp=realloc(*pppcLst,sizeof(char*)*(i+2ul));
+            if(!ppcTmp) {
+                fprintf(stderr,"Unable to allocate host memory\n");
+                perror("");
+                _ghf_freeHostZ(pcSource);
+                return(_GHM_HOST_MEMALLOC_ERROR);
+            }
+            *pppcLst=ppcTmp;
+            (*pppcLst)[i]=pcSource;
+            (*pppcLst)[i+1]=NULL;
+        }
+        return(_GHM_OK);
+    } else return(_GHM_NULL_POINTER_RECEIVED_ERROR);
+}
+
+void _ghf_wipeCharPtrList(char*** const pppcLst) {
+    if(pppcLst && *pppcLst) {
+        uint64_t i=0ul;
+        for(i=0ul; (*pppcLst)[i]; i++) _ghf_freeHostZ(&(*pppcLst)[i]);
+    }
+    _ghf_freeHostZ(pppcLst);
+    return;
+}

src/oclh_h_host_srr.c

@@ -0,0 +1,263 @@
+/*
+ * oclh_h_srr.c
+ *      Author: havock
+ */
+#include <stdio.h>
+#include <oclh_h_settings.h>
+#include <oclh_h_externals.h>
+#include <oclh_h_srr.h>
+
+#define _OCLH_OCL_HOST_ALGORITHMS_ 1
+#include <oclh_hd_srr.clh>
+#undef _OCLH_OCL_HOST_ALGORITHMS_
+
+_GHM_BASETYPE_HD_IMPLEMENTATION(i8)
+_GHM_BASETYPE_HD_IMPLEMENTATION(u8)
+_GHM_BASETYPE_HD_IMPLEMENTATION(i16)
+_GHM_BASETYPE_HD_IMPLEMENTATION(u16)
+_GHM_BASETYPE_HD_IMPLEMENTATION(i32)
+_GHM_BASETYPE_HD_IMPLEMENTATION(u32)
+_GHM_BASETYPE_HD_IMPLEMENTATION(i64)
+_GHM_BASETYPE_HD_IMPLEMENTATION(u64)
+_GHM_BASETYPE_HD_IMPLEMENTATION(f32)
+_GHM_BASETYPE_HD_IMPLEMENTATION(f64)
+
+/*
+ * host-device array of unsigned 8-bit values
+ * can be used as any data, considering byte alignment
+ */
+_GHT_HDvoid _ghf_declHD_void(void) {
+    const _GHT_HDvoid hdv={ .hPtr=NULL, .dPtr=NULL, .u64NofBytes=0ul };
+    return(hdv);
+}
+
+int32_t _ghf_wdcAllocHD_void(_GHT_WRKSET wSet,_GHT_HDvoid* const pHdv,
+                             const uint64_t u64NofBytes) {
+    if(!pHdv) return(_GHM_NULL_POINTER_RECEIVED_ERROR);
+    _ghf_wipeHD_void(wSet,pHdv);
+    pHdv->hPtr=_ghf_wdcAllocHostBuf(wSet, u64NofBytes);
+    if(!pHdv->hPtr) {
+        _ghf_wipeHD_void(wSet, pHdv); return(_GHM_HOST_MEMALLOC_ERROR);
+    }
+    pHdv->dPtr=_ghf_wdcAllocDevBuf(wSet,CL_MEM_READ_WRITE,u64NofBytes,NULL);
+    if(!pHdv->dPtr) {
+        _ghf_wipeHD_void(wSet,pHdv); return(_GHM_DEVICE_MEMALLOC_ERROR);
+    }
+    pHdv->u64NofBytes=u64NofBytes;
+    return(_GHM_OK);
+}
+
+int32_t _ghf_wdcReAllocHD_void(_GHT_WRKSET wSet, const _GHE_SYNC_TYPE SyncType,
+                               _GHT_HDvoid* const pHdv,
+                               const uint64_t u64NofBytes) {
+    uint8_t* hTmpPtr=NULL;
+    if(!pHdv) return(_GHM_NULL_POINTER_RECEIVED_ERROR);
+    hTmpPtr=realloc(pHdv->hPtr, u64NofBytes);
+    if(!hTmpPtr) {
+        if(wSet.pvDat) wSet.pfnDatCleaner(wSet);
+        _ghf_wipeHD_void(wSet,pHdv); return(_GHM_HOST_MEMALLOC_ERROR);
+    }
+    pHdv->hPtr=hTmpPtr;
+    _ghf_freeDevZ(wSet, &pHdv->dPtr);
+    pHdv->dPtr=_ghf_wdcAllocDevBuf(wSet, CL_MEM_READ_WRITE, u64NofBytes, NULL);
+    if(!pHdv->dPtr) {
+        _ghf_wipeHD_void(wSet,pHdv); return(_GHM_DEVICE_MEMALLOC_ERROR);
+    }
+    pHdv->u64NofBytes=u64NofBytes;
+    { int32_t err=0; 
+      if((err=_ghf_wdcSyncHD_void(wSet,SyncType,pHdv))) return(err); }
+    return(_GHM_OK);
+}
+
+int32_t _ghf_wdcSyncHD_void(_GHT_WRKSET wSet, const _GHE_SYNC_TYPE SyncType,
+                            _GHT_HDvoid* const pHdv) {
+    int32_t err=0;
+    if(!pHdv) return(_GHM_NULL_POINTER_RECEIVED_ERROR);
+    switch(SyncType) {
+        case _GHE_SYNC_HOST_TO_DEV:
+            {
+                char pcLogMsg[_GHM_MAXLEN_OF_LOGMSG];
+                snprintf(pcLogMsg, _GHM_MAXLEN_OF_LOGMSG,
+                         "%s/clEnqueueWriteBuffer", __func__);
+                wSet.APIErr=clEnqueueWriteBuffer(wSet.Queue, pHdv->dPtr,
+                                                 CL_TRUE, 0ul,
+                                                 pHdv->u64NofBytes,
+                                                 pHdv->hPtr, 0u, NULL, NULL);
+                if((err=_ghf_wdcChkWS_APIErr(wSet,pcLogMsg,_GHM_FL))) {
+                    _ghf_wipeHD_void(wSet, pHdv); return(err);
+                }
+            }
+            break;
+        case _GHE_SYNC_DEV_TO_HOST:
+            {
+                char pcLogMsg[_GHM_MAXLEN_OF_LOGMSG];
+                snprintf(pcLogMsg, _GHM_MAXLEN_OF_LOGMSG,
+                         "%s/clEnqueueReadBuffer", __func__);
+                wSet.APIErr=clEnqueueReadBuffer(wSet.Queue, pHdv->dPtr, CL_TRUE,
+                                                0ul, pHdv->u64NofBytes,
+                                                pHdv->hPtr, 0u, NULL, NULL);
+                if((err=_ghf_wdcChkWS_APIErr(wSet,pcLogMsg,_GHM_FL))) {
+                    _ghf_wipeHD_void(wSet, pHdv); return(err);
+                }
+            }
+            break;
+        default:
+            return(_GHM_UNKNOWN_SYNC_TYPE_ERROR);
+            break;
+    }
+    return(_GHM_OK);
+}
+
+int32_t _ghf_wdcFillHD_void(_GHT_WRKSET wSet, _GHT_HDvoid* const pHdv,
+                            const uint8_t u8Val) {
+    int32_t err=0;
+    uint64_t i=0ul;
+    if(!pHdv) return(_GHM_NULL_POINTER_RECEIVED_ERROR);
+    for(i=0ul; i<pHdv->u64NofBytes; i++) pHdv->hPtr[i]=u8Val;
+    if((err=_ghf_wdcSyncHD_void(wSet,_GHE_SYNC_HOST_TO_DEV,pHdv))) return(err);
+    return(_GHM_OK);
+}
+
+int32_t _ghf_wipeHD_void(_GHT_WRKSET wSet,_GHT_HDvoid* const pHdv) {
+    if(pHdv) {
+        _ghf_freeDevZ(wSet, &pHdv->dPtr);
+        _ghf_freeHostZ(&pHdv->hPtr);
+        pHdv->u64NofBytes=0ul;
+    }
+    return(_GHM_OK);
+}
+
+
+/*
+ * host-device array of float 32-bit vectors
+ */
+_GHT_HDVECS_F32 _ghf_declHDV_f32(void) {
+    const _GHT_HDVECS_F32 pVhdf32={
+            .hdf32V=_ghf_declHD_f32(),
+            .u64Dim=0ul, .u64SemDim=0ul
+    };
+    return(pVhdf32);
+}
+
+int32_t _ghf_wdcAllocHDV_f32(_GHT_WRKSET wSet, _GHT_HDVECS_F32* const pVhdf32,
+                             const uint64_t u64Dim, const uint64_t u64SemDim,
+                             const uint64_t u64NofVecs) {
+    int32_t err=0;
+    if(!pVhdf32) return(_GHM_NULL_POINTER_RECEIVED_ERROR);
+    _ghf_wipeHDV_f32(wSet,pVhdf32);
+    if((u64Dim<u64SemDim) || (u64Dim%u64SemDim)) return(4); /* TODO: make error codes */
+    if((err=_ghf_wdcAllocHD_f32(wSet, &pVhdf32->hdf32V, u64Dim*u64NofVecs))) {
+        _ghf_wipeHDV_f32(wSet,pVhdf32); return(err);
+    }
+    pVhdf32->u64Dim=u64Dim;
+    pVhdf32->u64SemDim=u64SemDim;
+    return(_GHM_OK);
+}
+
+int32_t _ghf_wdcSyncHDV_f32(_GHT_WRKSET wSet, const _GHE_SYNC_TYPE SyncType,
+                            _GHT_HDVECS_F32* const pVhdf32) {
+    int32_t err=0;
+    if(!pVhdf32) return(_GHM_NULL_POINTER_RECEIVED_ERROR);
+    if((err=_ghf_wdcSyncHD_f32(wSet,SyncType,&pVhdf32->hdf32V))) return(err); /* TODO: make error codes */
+    return(_GHM_OK);
+}
+
+int32_t _ghf_wdcFillHDV_f32(_GHT_WRKSET wSet, _GHT_HDVECS_F32* const pVhdf32,
+                            const flt32_t f32Val) {
+    int32_t err=0;
+    uint64_t i=0ul;
+    if(!pVhdf32) return(_GHM_NULL_POINTER_RECEIVED_ERROR);
+    for(i=0ul; i<pVhdf32->hdf32V.u64NofVals; i++)
+        pVhdf32->hdf32V.hPtr[i]=f32Val;
+    if((err=_ghf_wdcSyncHDV_f32(wSet,_GHE_SYNC_HOST_TO_DEV,pVhdf32)))
+        return(err);
+    return(_GHM_OK);
+}
+
+int32_t _ghf_isHDV_f32_Valid(const _GHT_HDVECS_F32 Vhdf32) {
+    if(Vhdf32.u64Dim &&
+       Vhdf32.u64SemDim && 
+       _ghf_isHD_f32_Valid(Vhdf32.hdf32V)) {
+        if((Vhdf32.u64Dim < Vhdf32.u64SemDim ) ||
+           (Vhdf32.u64Dim % Vhdf32.u64SemDim ) ||
+           (Vhdf32.hdf32V.u64NofVals % Vhdf32.u64Dim )) return(0);
+        return(1);
+    }
+    return(0); /* Внимание! 0 -- значит набор векторов невалиден */
+}
+
+uint64_t _ghf_getNofVecsHDV_f32(const _GHT_HDVECS_F32 Vhdf32) {
+    if(!_ghf_isHDV_f32_Valid(Vhdf32))
+        return((uint64_t) -1);/* max unsigned long int */
+    return(Vhdf32.hdf32V.u64NofVals/Vhdf32.u64Dim);
+}
+
+int32_t _ghf_saveHDV_f32_ToTxt(const _GHT_HDVECS_F32 vfVecs,
+                               const char* const pcFileName,
+                               const int32_t i32UseSemDimFlag) {
+    if(vfVecs.u64Dim &&            vfVecs.u64SemDim &&
+       vfVecs.hdf32V.u64NofVals && vfVecs.hdf32V.hPtr) {
+        FILE* pFile=NULL;
+        uint64_t i=0ul;
+        const uint64_t u64Dim=(i32UseSemDimFlag)?vfVecs.u64SemDim:vfVecs.u64Dim,
+                       u64NofV=vfVecs.hdf32V.u64NofVals/u64Dim;
+        pFile=fopen(pcFileName,"w");
+        if(!pFile) return(_GHM_OPEN_FILE_ERROR);
+        fprintf(pFile,"seqn\t");
+        for(i=0ul; i<u64Dim; i++) {
+            if(i==(u64Dim-1)) fprintf(pFile,"c%lu\n",i);
+            else fprintf(pFile,"c%lu\t",i);
+        }
+        for(i=0ul; i<u64NofV; i++) {
+            uint64_t j=0ul;
+            flt32_t* pf32CurVec=&vfVecs.hdf32V.hPtr[i*u64Dim];
+            fprintf(pFile,"%f\t",(flt32_t) i);
+            for(j=0ul; j<u64Dim; j++) {
+                if(j==(u64Dim-1)) fprintf(pFile,"%f\n",pf32CurVec[j]);
+                else fprintf(pFile,"%f\t",pf32CurVec[j]);
+            }
+        }
+        fclose(pFile);
+        return(_GHM_OK);
+    }
+    return(1); /* TODO: make error codes */
+}
+int32_t _ghf_saveHDV_f32_ToTxtByHD_u64_Idxs(const _GHT_HDVECS_F32 vfVecs,
+                                            const _GHT_HDu64 hdu64IV,
+                                            const char* const pcFileName,
+                                            const int32_t i32UseSemDimFlag) {
+    if(vfVecs.u64Dim &&            vfVecs.u64SemDim &&
+       vfVecs.hdf32V.u64NofVals && vfVecs.hdf32V.hPtr &&
+       hdu64IV.u64NofVals &&       hdu64IV.hPtr) {
+        const uint64_t u64Dim=(i32UseSemDimFlag)?vfVecs.u64SemDim:vfVecs.u64Dim;
+        FILE* pFile=NULL;
+        uint64_t i=0ul;
+        pFile=fopen(pcFileName, "w");
+        if(!pFile) return(_GHM_OPEN_FILE_ERROR);
+        fprintf(pFile, "seqn\t");
+        for(i=0ul; i<u64Dim; i++) {
+            if(i==(u64Dim-1)) fprintf(pFile,"c%lu\n",i);
+            else fprintf(pFile,"c%lu\t",i);
+        }
+        for(i=0ul; i<hdu64IV.u64NofVals; i++) {
+            uint64_t j=0ul;
+            flt32_t* pf32CurVec=&vfVecs.hdf32V.hPtr[hdu64IV.hPtr[i]*u64Dim];
+            fprintf(pFile,"%f\t",(flt32_t) i);
+            for(j=0ul; j<u64Dim; j++) {
+                if(j==(u64Dim-1)) fprintf(pFile,"%f\n",pf32CurVec[j]);
+                else fprintf(pFile,"%f\t",pf32CurVec[j]);
+            }
+        }
+        fclose(pFile);
+        return(_GHM_OK);
+    }
+    return(1); /* TODO: make error codes */
+}
+int32_t _ghf_wipeHDV_f32(_GHT_WRKSET wSet, _GHT_HDVECS_F32* const pVhdf32) {
+    if(pVhdf32) {
+        _ghf_wipeHD_f32(wSet,&pVhdf32->hdf32V);
+        pVhdf32->u64Dim=0ul;
+        pVhdf32->u64SemDim=0ul;
+    }
+    return(_GHM_OK);
+}

src/oclh_h_internals.c

@@ -0,0 +1,57 @@
+/*
+ * oclh_h_internals.c
+ *      Author: havock
+ */
+#include <stdlib.h>
+#include <string.h>
+#include <ctype.h>
+#include <oclh_h_internals.h>
+#include <oclh_h_errors.h>
+
+
+uint64_t __ghf_removePreNPostSpacesFromCharPtr(char* const pcStr) {
+    if(!pcStr || !pcStr[0]) return(0ul);
+    uint64_t u64NofRemovedSpaces=0ul;
+    { /* preSpaces remove */
+        uint64_t i=0ul;
+        while(pcStr[i] && isspace(pcStr[i])) i++;
+        if(i && pcStr[i]!='\0') {
+            uint64_t j=i, k=0ul;
+            while(pcStr[j]) pcStr[k++]=pcStr[j++];
+            pcStr[k]='\0';
+            u64NofRemovedSpaces+=i;
+        }
+    }
+    { /*postSpaces remove */
+        uint64_t i=strlen(pcStr);
+        while(isspace(pcStr[i]) && i) {
+            pcStr[i]='\0'; i--; u64NofRemovedSpaces++;
+        }
+        if(!i && isspace(pcStr[i])) { pcStr[i]='\0'; u64NofRemovedSpaces++; }
+    }
+    return(u64NofRemovedSpaces);
+}
+
+uint64_t __ghf_replaceSpacesWithUnderscoreInCharPtr(char* const pcStr) {
+    if(pcStr) {
+        uint64_t u64NofReplacedSpaces=0ul, i=0ul;
+        while(pcStr[i]) {
+            if(isspace(pcStr[i])) { pcStr[i]='_'; u64NofReplacedSpaces++; }
+            i++;
+        }
+        return(u64NofReplacedSpaces);
+    } else return(0ul);
+}
+
+int32_t  __ghf_concatHeapStrAndCharPtr(char** ppcDst, const char* const pcStr) {
+    if(ppcDst) {
+        char* pcTmp=NULL;
+        pcTmp=realloc(*ppcDst,
+                    (*ppcDst?strlen(*ppcDst):0ul)+
+                    (pcStr?strlen(pcStr):0ul)+2ul);
+        if(!pcTmp) return(_GHM_HOST_MEMALLOC_ERROR);
+        *ppcDst=pcTmp;
+        strncat(*ppcDst," ",1ul); strncat(*ppcDst,pcStr,strlen(pcStr));
+        return(_GHM_OK);
+    } else return(_GHM_NULL_POINTER_RECEIVED_ERROR);
+}

src/oclh_h_ws_base_log.c

@@ -0,0 +1,324 @@
+/*
+ * oclh_h_ws_base_log.c
+ *      Author: havock
+ */
+#include <time.h>
+#include <string.h>
+#include <oclh_h_settings.h>
+#include <oclh_h_base_clapi_strings.h>
+#include <oclh_h_ws_base_log.h>
+#include <oclh_hd_std_types.clh>
+#include <oclh_h_internals.h>
+
+int32_t __ghf_logWS_MsgIgnoringLock(const _GHT_WRKSET wSet,
+                                    const char* const pcLogMsg) {
+    time_t timeRaw;
+    struct tm* stTime;
+    time(&timeRaw);
+    stTime=localtime(&timeRaw);
+    fprintf(wSet.Log.pfOut,"%04d-%02d-%02d %02d:%02d:%02d ws_0x%04lx %s\n",
+            stTime->tm_year+1900,stTime->tm_mon+1,stTime->tm_mday,
+            stTime->tm_hour,stTime->tm_min,stTime->tm_sec,
+            __GHM_U64STRIPTO2B(wSet.pwSetAddr),pcLogMsg);
+#ifdef _GHM_FLUSH_LOGS
+    fflush(wSet.Log.pfOut);
+#endif /* _GHM_FLUSH_LOGS */
+    return(_GHM_OK);
+}
+
+int32_t _ghf_logWS_Msg(const _GHT_WRKSET wSet, const char* const pcLogMsg) {
+    if(_ghf_isLogValid(wSet.Log)) {
+        pthread_mutex_lock(wSet.Log.pMtx);
+            __ghf_logWS_MsgIgnoringLock(wSet, pcLogMsg);
+        pthread_mutex_unlock(wSet.Log.pMtx);
+        return(_GHM_OK);
+    }
+    return(_GHM_NULL_POINTER_RECEIVED_ERROR);
+}
+
+int32_t _ghf_logWS_Hdr(const _GHT_WRKSET wSet,const char* const pcLogHdr) {
+    char pcLogHdrTop[_GHM_MAXLEN_OF_LOGMSG];
+    char pcLogHdrBottom[_GHM_MAXLEN_OF_LOGMSG];
+    const size_t len=strlen(pcLogHdr);
+    size_t i=0ul;
+    for(i=0ul;i<len && i<_GHM_MAXLEN_OF_LOGMSG-2ul;i++) {
+        pcLogHdrTop[i]='_'; pcLogHdrBottom[i]='~';
+    }
+    pcLogHdrTop[i]=pcLogHdrBottom[i]='\0';
+    pthread_mutex_lock(wSet.Log.pMtx);
+        __ghf_logWS_MsgIgnoringLock(wSet,pcLogHdrTop);
+        __ghf_logWS_MsgIgnoringLock(wSet,pcLogHdr);
+        __ghf_logWS_MsgIgnoringLock(wSet,pcLogHdrBottom);
+    pthread_mutex_unlock(wSet.Log.pMtx);
+    return(_GHM_OK);
+}
+
+int32_t _ghf_logWS_Delim(const _GHT_WRKSET wSet) {
+    pthread_mutex_lock(wSet.Log.pMtx);
+        __ghf_logWS_MsgIgnoringLock(wSet,
+             "_______________________________________________________________");
+        __ghf_logWS_MsgIgnoringLock(wSet,
+             "~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~");
+    pthread_mutex_unlock(wSet.Log.pMtx);
+    return(_GHM_OK);
+}
+
+int32_t _ghf_logWS_Raw(const _GHT_WRKSET wSet, const char* const pcLogMessage) {
+    pthread_mutex_lock(wSet.Log.pMtx);
+        fprintf(wSet.Log.pfOut,"%s\n",pcLogMessage);
+#ifdef _GHM_FLUSH_LOGS
+        fflush(wSet.Log.pfOut);
+#endif /* _GHM_FLUSH_LOGS */
+    pthread_mutex_unlock(wSet.Log.pMtx);
+    return(_GHM_OK);
+}
+
+int32_t _ghf_logWS_APIErr(const _GHT_WRKSET wSet, const char* const pcAPICall) {
+    char pcLogMsg[_GHM_MAXLEN_OF_LOGMSG];
+    snprintf(pcLogMsg,_GHM_MAXLEN_OF_LOGMSG,
+             "oclerr: %s returned error %d - %s",
+             pcAPICall,wSet.APIErr, _ghf_CLAPIErrString(wSet.APIErr));
+    _ghf_logWS_Msg(wSet,pcLogMsg);
+    return(_GHM_OK);
+}
+
+int32_t _ghf_logWS_UsedMem(const _GHT_WRKSET wSet) {
+    if(wSet.pszMemBytes) {
+        char pcLogMsg[_GHM_MAXLEN_OF_LOGMSG];
+        snprintf(pcLogMsg,_GHM_MAXLEN_OF_LOGMSG,
+                 "Mem used: %lu B/%.2f KB/%.2f MB/%.2f GB",
+                 (*wSet.pszMemBytes),
+                 ((flt32_t) *wSet.pszMemBytes)/1024e0f,
+                 ((flt32_t) *wSet.pszMemBytes)/1024e0f/1024e0f,
+                 ((flt32_t) *wSet.pszMemBytes)/1024e0f/1024e0f/1024e0f);
+        _ghf_logWS_Msg(wSet,pcLogMsg);
+    }
+    return(_GHM_OK);
+}
+
+int32_t _ghf_logWS_DevInf_bool(      _GHT_WRKSET    wSet,
+                               const cl_device_id   clDev,
+                               const cl_device_info InfoVal,
+                               const char*    const pcPrefix,
+                               const char*    const pcName,
+                               const char*    const pcTrue,
+                               const char*    const pcFalse) {
+    char pcLogMsg[_GHM_MAXLEN_OF_LOGMSG];
+    size_t szRes=0ul;
+    cl_bool clbInf=0u;
+    wSet.APIErr=clGetDeviceInfo(clDev,InfoVal,sizeof(size_t),&clbInf,&szRes);
+    if(wSet.APIErr)
+        snprintf(pcLogMsg,_GHM_MAXLEN_OF_LOGMSG,
+                 "%s %s Undefined! clGetDeviceInfo returned oclerr %s(%d)",
+                 pcPrefix?pcPrefix:"",pcName,
+                 _ghf_CLAPIErrString(wSet.APIErr),wSet.APIErr);
+    else snprintf(pcLogMsg,_GHM_MAXLEN_OF_LOGMSG,
+                  "%s %s %s",pcPrefix?pcPrefix:"",pcName,
+                  clbInf?pcTrue:pcFalse);
+    _ghf_logWS_Msg(wSet,pcLogMsg);
+    return(_GHM_OK);
+}
+
+int32_t _ghf_logWS_DevInf_charptr(      _GHT_WRKSET    wSet,
+                                  const cl_device_id   clDev,
+                                  const cl_device_info InfoVal,
+                                  const char*    const pcPrefix,
+                                  const char*    const pcName,
+                                  const char*    const pcPostfix) {
+    char pcLogMsg[_GHM_MAXLEN_OF_LOGMSG], *pcInf=NULL;
+    size_t szRes=0ul;
+    wSet.APIErr=clGetDeviceInfo(clDev, InfoVal, 0ul, NULL, &szRes);
+    if(wSet.APIErr) {
+        snprintf(pcLogMsg, _GHM_MAXLEN_OF_LOGMSG,
+                 "%s %s Undefined! clGetDeviceInfo returned oclerr %s(%d)",
+                 pcPrefix?pcPrefix:"", pcName,
+                 _ghf_CLAPIErrString(wSet.APIErr), wSet.APIErr);
+    } else {
+        if(!szRes) {
+            if(pcPrefix) snprintf(pcLogMsg, _GHM_MAXLEN_OF_LOGMSG,
+                                  "%s %s empty value", pcPrefix, pcName);
+            else snprintf(pcLogMsg, _GHM_MAXLEN_OF_LOGMSG,
+                          "%s empty value", pcName);
+            return(_GHM_OK);
+        }
+        pcInf=malloc(szRes);
+        if(pcInf) {
+            wSet.APIErr=clGetDeviceInfo(clDev, InfoVal, szRes, pcInf, NULL);
+            if(wSet.APIErr)
+                snprintf(pcLogMsg, _GHM_MAXLEN_OF_LOGMSG,
+                         "%s %s Undefined! "
+                         "clGetDeviceInfo returned oclerr %s(%d)",
+                         pcPrefix?pcPrefix:"", pcName,
+                         _ghf_CLAPIErrString(wSet.APIErr), wSet.APIErr);
+            else {
+                pcInf[szRes-1]='\0';
+                __ghf_removePreNPostSpacesFromCharPtr(pcInf);
+                if(pcPrefix)
+                    snprintf(pcLogMsg, _GHM_MAXLEN_OF_LOGMSG, "%s %s %s %s",
+                             pcPrefix, pcName, pcInf, pcPostfix?pcPostfix:"" );
+                else snprintf(pcLogMsg, _GHM_MAXLEN_OF_LOGMSG,
+                              "%s %s", pcName, pcInf);
+            }
+            if(pcInf) { free(pcInf); pcInf=NULL; }
+        } else snprintf(pcLogMsg, _GHM_MAXLEN_OF_LOGMSG,
+                        "%s %s Undefined! Host memory allocation failed",
+                        pcPrefix?pcPrefix:"", pcName);
+    }
+    _ghf_logWS_Msg(wSet, pcLogMsg);
+    return(_GHM_OK);
+}
+
+int32_t _ghf_logWS_DevInf_cluint(      _GHT_WRKSET    wSet,
+                                 const cl_device_id   clDev,
+                                 const cl_device_info InfoVal,
+                                 const char*    const pcPrefix,
+                                 const char*    const pcName,
+                                 const char*    const pcPostfix) {
+    char pcLogMsg[_GHM_MAXLEN_OF_LOGMSG];
+    size_t szRes=0ul;
+    cl_uint cluInf=0u;
+    wSet.APIErr=clGetDeviceInfo(clDev,InfoVal,sizeof(cl_uint),&cluInf,&szRes);
+    if(wSet.APIErr)
+        snprintf(pcLogMsg ,_GHM_MAXLEN_OF_LOGMSG,
+                 "%s %s Undefined! clGetDeviceInfo returned oclerr %s(%d)",
+                 pcPrefix?pcPrefix:"", pcName,
+                 _ghf_CLAPIErrString(wSet.APIErr), wSet.APIErr);
+    else {
+        if(!strncmp("bytes",pcPostfix?pcPostfix:"",_GHM_MAXLEN_OF_LOGMSG)) {
+            if(pcPrefix)
+                snprintf(pcLogMsg,_GHM_MAXLEN_OF_LOGMSG,
+                         "%s %s %.2f GiB/%.2f MiB/%.2f KiB/%u bytes",
+                         pcPrefix,pcName,
+                         ((float)cluInf)/1024e0f/1024e0f/1024e0f,
+                         ((float)cluInf)/1024e0f/1024e0f,
+                         ((float)cluInf)/1024e0f,cluInf);
+            else snprintf(pcLogMsg,_GHM_MAXLEN_OF_LOGMSG,
+                         "%s %.2f GiB/%.2f MiB/%.2f KiB/%u bytes",
+                         pcName,((float)cluInf)/1024e0f/1024e0f/1024e0f,
+                         ((float)cluInf)/1024e0f/1024e0f,
+                         ((float)cluInf)/1024e0f,cluInf);
+        } else {
+            if(pcPrefix)
+                snprintf(pcLogMsg, _GHM_MAXLEN_OF_LOGMSG,
+                         "%s %s %u %s", pcPrefix, pcName, cluInf,
+                         pcPostfix?pcPostfix:"");
+            else snprintf(pcLogMsg, _GHM_MAXLEN_OF_LOGMSG,
+                          "%s %u %s", pcName, cluInf, pcPostfix?pcPostfix:"");
+        }
+    }
+    _ghf_logWS_Msg(wSet, pcLogMsg);
+    return(_GHM_OK);
+}
+
+int32_t _ghf_logWS_DevInf_clulong(      _GHT_WRKSET    wSet,
+                                  const cl_device_id   clDev,
+                                  const cl_device_info InfoVal,
+                                  const char*    const pcPrefix,
+                                  const char*    const pcName,
+                                  const char*    const pcPostfix) {
+    char pcLogMsg[_GHM_MAXLEN_OF_LOGMSG];
+    size_t szResSize=0ul;
+    cl_ulong clulInf=0ul;
+    wSet.APIErr=
+        clGetDeviceInfo(clDev, InfoVal, sizeof(cl_ulong), &clulInf, &szResSize);
+    if(wSet.APIErr)
+        snprintf(pcLogMsg, _GHM_MAXLEN_OF_LOGMSG,
+                 "%s %s Undefined! clGetDeviceInfo returned oclerr %s(%d)",
+                 pcPrefix?pcPrefix:"", pcName,
+                 _ghf_CLAPIErrString(wSet.APIErr), wSet.APIErr);
+    else {
+        if(!strncmp("bytes",pcPostfix?pcPostfix:"",_GHM_MAXLEN_OF_LOGMSG)) {
+            if(pcPrefix)
+                snprintf(pcLogMsg, _GHM_MAXLEN_OF_LOGMSG,
+                         "%s %s %.2f GiB/%.2f MiB/%.2f KiB/%lu bytes",
+                         pcPrefix, pcName,
+                         ((float) clulInf)/1024e0f/1024e0f/1024e0f,
+                         ((float) clulInf)/1024e0f/1024e0f,
+                         ((float) clulInf)/1024e0f,clulInf);
+            else snprintf(pcLogMsg, _GHM_MAXLEN_OF_LOGMSG,
+                          "%s %.2f GiB/%.2f MiB/%.2f KiB/%lu bytes",
+                          pcName, ((float) clulInf)/1024e0f/1024e0f/1024e0f,
+                          ((float) clulInf)/1024e0f/1024e0f,
+                          ((float) clulInf)/1024e0f,clulInf);
+        } else {
+            if(pcPrefix)
+                snprintf(pcLogMsg, _GHM_MAXLEN_OF_LOGMSG,
+                         "%s %s %lu %s", pcPrefix, pcName, clulInf,
+                         pcPostfix?pcPostfix:"");
+            else snprintf(pcLogMsg, _GHM_MAXLEN_OF_LOGMSG,
+                          "%s %lu %s", pcName, clulInf, pcPostfix?pcPostfix:"");
+        }
+    }
+    _ghf_logWS_Msg(wSet, pcLogMsg);
+    return(_GHM_OK);
+}
+
+int32_t _ghf_logWS_DevInf_clulong_as_hex(      _GHT_WRKSET    wSet,
+                                         const cl_device_id   clDev,
+                                         const cl_device_info InfoVal,
+                                         const char*    const pcPrefix,
+                                         const char*    const pcName,
+                                         const char*    const pcPostfix) {
+    char pcLogMsg[_GHM_MAXLEN_OF_LOGMSG];
+    size_t szResSize=0ul;
+    cl_ulong clulInf=0ul;
+    wSet.APIErr=
+        clGetDeviceInfo(clDev, InfoVal, sizeof(cl_ulong), &clulInf, &szResSize);
+    if(wSet.APIErr)
+        snprintf(pcLogMsg, _GHM_MAXLEN_OF_LOGMSG,
+                 "%s %s Undefined! clGetDeviceInfo returned oclerr %s(%d)",
+                 pcPrefix?pcPrefix:"", pcName,
+                 _ghf_CLAPIErrString(wSet.APIErr), wSet.APIErr);
+    else {
+        if(pcPrefix)
+            snprintf(pcLogMsg, _GHM_MAXLEN_OF_LOGMSG,
+                         "%s %s 0x%lx %s", pcPrefix, pcName, clulInf,
+                         pcPostfix?pcPostfix:"");
+            else snprintf(pcLogMsg, _GHM_MAXLEN_OF_LOGMSG,
+                          "%s 0x%lx %s", pcName, clulInf,
+                          pcPostfix?pcPostfix:"");
+    }
+    _ghf_logWS_Msg(wSet,pcLogMsg);
+    return(_GHM_OK);
+}
+
+int32_t _ghf_logWS_DevInf_size(      _GHT_WRKSET    wSet,
+                               const cl_device_id   clDev,
+                               const cl_device_info InfoVal,
+                               const char*    const pcPrefix,
+                               const char*    const pcName,
+                               const char*    const pcPostfix) {
+    char pcLogMsg[_GHM_MAXLEN_OF_LOGMSG];
+    size_t szRes=0ul, szInf=0ul;
+    wSet.APIErr=clGetDeviceInfo(clDev, InfoVal, sizeof(size_t), &szInf, &szRes);
+    if(wSet.APIErr)
+        snprintf(pcLogMsg, _GHM_MAXLEN_OF_LOGMSG,
+                 "%s %s Undefined! clGetDeviceInfo returned oclerr %s(%d)",
+                 pcPrefix?pcPrefix:"", pcName,
+                 _ghf_CLAPIErrString(wSet.APIErr), wSet.APIErr);
+    else {
+        if(!strncmp("bytes",pcPostfix?pcPostfix:"",_GHM_MAXLEN_OF_LOGMSG)) {
+            if(pcPrefix)
+                snprintf(pcLogMsg, _GHM_MAXLEN_OF_LOGMSG,
+                         "%s %s %.2f GiB/%.2f MiB/%.2f KiB/%lu bytes",
+                         pcPrefix, pcName,
+                         ((float) szInf)/1024e0f/1024e0f/1024e0f,
+                         ((float) szInf)/1024e0f/1024e0f,
+                         ((float) szInf)/1024e0f,szInf);
+            else snprintf(pcLogMsg, _GHM_MAXLEN_OF_LOGMSG,
+                          "%s %.2f GiB/%.2f MiB/%.2f KiB/%lu bytes",
+                          pcName, ((float) szInf)/1024e0f/1024e0f/1024e0f,
+                          ((float) szInf)/1024e0f/1024e0f,
+                          ((float) szInf)/1024e0f,szInf);
+        } else {
+            if(pcPrefix)
+                snprintf(pcLogMsg, _GHM_MAXLEN_OF_LOGMSG,
+                         "%s %s %lu %s", pcPrefix, pcName, szInf,
+                         pcPostfix?pcPostfix:"");
+            else snprintf(pcLogMsg, _GHM_MAXLEN_OF_LOGMSG,
+                          "%s %lu %s", pcName, szInf, pcPostfix?pcPostfix:"");
+        }
+    }
+    _ghf_logWS_Msg(wSet,pcLogMsg);
+    return(_GHM_OK);
+}