/*
 * oclh_d_mem_alloc.clh
 *      Author: havock
 */
//#define _OCLH_OCL_COMPILER_
#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 {          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__NULL(__GDM_NULL_MAC_Region) \
    ((uint64_t) __GD__ ## __GDM_NULL_MAC_Region ## _Heap_u8Ptr)
#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_is_ptr_NULL(__GDM_ISNULL_MAC_Region,__GDM_ISNULL_MAC_Ptr) \
   (((uint64_t) (__GDM_ISNULL_MAC_Ptr))==_GDM__NULL(__GDM_ISNULL_MAC_Region))
#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 static uint8_t* malloc__private(__private uint8_t* const pu8HeapPtr,
                                              const uint64_t       u64Size);
__global  static uint8_t* malloc__global( __global  uint8_t* const pu8HeapPtr,
                                              const uint64_t       u64Size);
static void free__private(__private uint8_t* const pu8HeapPtr,
                   __private uint8_t**      ppu8Ptr);
static void free__global( __global  uint8_t* const pu8HeapPtr,
                   __global  uint8_t**      ppu8Ptr);
static int32_t Print__private_HeapState(__private const uint8_t
                                                            * const pu8HeapPtr);
static int32_t Print__global_HeapState( __global  const uint8_t
                                                            * const pu8HeapPtr);

__private static 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 static 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);
}
static 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;
}
static 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;
}
static 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);
}
static 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);
}

#endif /* _OCLH_OCL_COMPILER_ */
#endif /* OCLH_D_MEM_ALLOC_CLH_ */
