nm_data.h

Go to the documentation of this file.

/*
 * NewMadeleine
 * Copyright (C) 2015-2026 (see AUTHORS file)
 *
 * This program is free software; you can redistribute it and/or modify
 * it under the terms of the GNU General Public License as published by
 * the Free Software Foundation; either version 2 of the License, or (at
 * your option) any later version.
 *
 * This program is distributed in the hope that it will be useful, but
 * WITHOUT ANY WARRANTY; without even the implied warranty of
 * MERCHANTABILITY or FITNESS FOR A PARTICULA R PURPOSE.  See the GNU
 * General Public License for more details.
 */
 
#ifndef NM_DATA_H
#define NM_DATA_H
 
#include <assert.h>
#include <string.h>
 
#include <Padico/Puk.h>
 
#include <nm_config.h>
 
#ifdef NMAD_CUDA
#include <cuda.h>
#include <cuda_runtime.h>
#include <cuda_runtime_api.h>
#endif /* NMAD_CUDA */
 
#ifdef NMAD_HIP
#include <hip/hip_runtime.h>
#include <hip/hip_runtime_api.h>
#endif /* NMAD_HIP */
 
/* ** Data descriptor ************************************** */
 
#define _NM_DATA_CONTENT_SIZE 64
 
/* forward declaration so that operators declared inside struct nm_data_s may
 * take struct nm_data_s* as parameters
 */
struct nm_data_s;
 
typedef void (*nm_data_apply_t)(void*ptr, nm_len_t len, void*_context);
 
struct nm_data_properties_s
{
  nm_len_t blocks; 
  nm_len_t size;   
  int is_contig;   
#ifdef NMAD_CUDA
  int is_cuda;     
#endif /* NMAD_CUDA */
#ifdef NMAD_HIP
  int is_hip;     
#endif /* NMAD_HIP */
};
 
typedef enum nm_data_op_kind_e
  {
    NM_DATA_OP_NONE = 0,        
    NM_DATA_OP_NOP,             
    NM_DATA_OP_COPY_FROM,       
    NM_DATA_OP_COPY_TO,         
    NM_DATA_OP_DYNAMIC,         
    NM_DATA_OP_GETPROPS,        
    NM_DATA_OP_CHUNK,           
    NM_DATA_OP_SLICER_COROUTINE 
  } nm_data_op_t;
 
struct nm_data_op_context_nop_s
{
  /* empty */
};
 
struct nm_data_op_context_copy_from_s
{
  void*__restrict__ p_dest_ptr;
  const struct nm_data_s*p_data;
};
 
struct nm_data_op_context_copy_to_s
{
  const void*__restrict__ p_src_ptr;
  const struct nm_data_s*p_data;
};
 
struct nm_data_op_context_dynamic_s
{
  nm_data_apply_t p_apply;
  void*p_apply_context;
};
 
struct nm_data_op_context_getprops_s
{
  void*p_blockend;  
  struct nm_data_properties_s props;
};
 
struct nm_data_op_context_chunk_s
{
  nm_len_t chunk_offset;    
  nm_len_t chunk_len;       
  nm_len_t done;            
  struct nm_data_op_s*p_op; 
};
 
struct nm_data_op_context_slicer_coroutine_s
{
  struct nm_data_slicer_s*p_slicer;
};
 
struct nm_data_op_s
{
  enum nm_data_op_kind_e kind;
  union
  {
    struct nm_data_op_context_nop_s nop;
    struct nm_data_op_context_copy_from_s copy_from;
    struct nm_data_op_context_copy_to_s copy_to;
    struct nm_data_op_context_dynamic_s dynamic;
    struct nm_data_op_context_getprops_s getprops;
    struct nm_data_op_context_chunk_s chunk;
    struct nm_data_op_context_slicer_coroutine_s slicer_coroutine;
  } context;
};
 
typedef void (*nm_data_traversal_t)(const void*_data_content, struct nm_data_op_s*p_op);
 
typedef void (*nm_data_properties_compute_t)(struct nm_data_s*p_data);
 
struct nm_data_ops_s
{
  nm_data_traversal_t          p_traversal;          
  nm_data_properties_compute_t p_properties_compute; 
};
 
void nm_data_default_properties_compute(struct nm_data_s*p_data);
 
struct nm_data_s
{
  struct nm_data_ops_s ops;              
  struct nm_data_properties_s props;     
  char _content[_NM_DATA_CONTENT_SIZE];  
};
 
static inline void nm_data_propertie_gpu_preinit(struct nm_data_properties_s*p_props);
 
static inline void nm_data_propertie_gpu_postinit(const struct nm_data_properties_s*p_props);
 
void nm_data_properties_gpu_fill(struct nm_data_properties_s*p_props, const void*p_ptr);
 
#define NM_DATA_TYPE(ENAME, CONTENT_TYPE, OPS)                          \
  __attribute__((unused))                                               \
  static inline void nm_data_##ENAME##_set(struct nm_data_s*p_data, CONTENT_TYPE value) \
  {                                                                     \
    p_data->ops = *(OPS);                                               \
    assert(p_data->ops.p_traversal != NULL);                            \
    if(p_data->ops.p_properties_compute == NULL)                        \
      {                                                                 \
        p_data->ops.p_properties_compute = nm_data_default_properties_compute; \
      }                                                                 \
    assert(sizeof(CONTENT_TYPE) <= _NM_DATA_CONTENT_SIZE);              \
    CONTENT_TYPE*p_content = (CONTENT_TYPE*)&p_data->_content[0];       \
    *p_content = value;                                                 \
    p_data->props.blocks = -1;                                          \
    nm_data_propertie_gpu_preinit(&p_data->props);                      \
    (*p_data->ops.p_properties_compute)(p_data);                        \
    nm_data_propertie_gpu_postinit(&p_data->props);                     \
  }                                                                     \
  __attribute__((unused))                                               \
  static inline CONTENT_TYPE*nm_data_##ENAME##_content(const struct nm_data_s*p_data) \
  {                                                                     \
    return (CONTENT_TYPE*)p_data->_content;                             \
  }
 
/* ** datav data (dynamic vector of nm_data)
 */
 
#define NM_DATAV_INIT_SIZE 4
 
struct nm_datav_s
{
  struct nm_data_s*p_data; 
  struct nm_data_s data[NM_DATAV_INIT_SIZE];  
  int n_data;             
  int allocated;          
  int commited;
};
 
static inline void nm_datav_init(struct nm_datav_s*p_datav);
static inline void nm_datav_destroy(struct nm_datav_s*p_datav);
static inline void nm_datav_add_chunk_data(struct nm_datav_s*p_datav, const struct nm_data_s*p_data);
static inline void nm_datav_add_chunk(struct nm_datav_s*p_datav, const void*ptr, nm_len_t len);
static inline nm_len_t nm_datav_size(struct nm_datav_s*p_datav);
static inline void nm_datav_uncommit(struct nm_datav_s*p_datav);
 
/* ** Built-in data ops ************************************ */
 
static inline void nm_data_op_apply(void*__restrict__ p_ptr, nm_len_t len, struct nm_data_op_s*__restrict__ p_op);
static inline void nm_data_op_apply_nop(void*__restrict__ p_ptr, nm_len_t len, struct nm_data_op_context_nop_s*__restrict__ p_nop);
static inline void nm_data_op_apply_copy_from(void*__restrict__ p_ptr, nm_len_t len, struct nm_data_op_context_copy_from_s*__restrict__ p_copy_from);
static inline void nm_data_op_apply_copy_to(void*__restrict__ p_ptr, nm_len_t len, struct nm_data_op_context_copy_to_s*__restrict__ p_copy_to);
static inline void nm_data_op_apply_dynamic(void*__restrict__ p_ptr, nm_len_t len, struct nm_data_op_context_dynamic_s*__restrict__ p_dynamic);
static inline void nm_data_op_apply_getprops(void*__restrict__ p_ptr, nm_len_t len, struct nm_data_op_context_getprops_s*__restrict__ p_getprops);
static inline void nm_data_op_apply_chunk(void*__restrict__ p_ptr, nm_len_t len, struct nm_data_op_context_chunk_s*__restrict__ p_chunk);
static inline void nm_data_op_apply_slicer_coroutine(void*__restrict__ p_ptr, nm_len_t len, struct nm_data_op_context_slicer_coroutine_s*__restrict__ p_slicer_coroutine);
 
#define NM_DATA_OP_APPLY_LOOP_REF(I, COUNT, PTR, LEN, P_OP)      \
  {                                                              \
    nm_len_t I;                                                  \
    for(I = 0; I < (COUNT); I++)                                 \
      {                                                          \
        nm_data_op_apply((PTR), (LEN), (P_OP));                  \
      }                                                          \
  }
 
#define NM_DATA_OP_APPLY_KIND_LEN(I, OP_KIND, START, COUNT, PTR, LEN, P_OP) \
  {                                                                     \
    nm_len_t I;                                                         \
    for(I = START; I < (START) + (COUNT); I++)                          \
      {                                                                 \
        nm_data_op_apply_##OP_KIND((PTR), (LEN), &(P_OP)->context.OP_KIND); \
      }                                                                 \
  }
 
#define NM_DATA_OP_APPLY_KIND_VECT(I, OP_KIND, START, COUNT, PTR, LEN, P_OP) \
  {                                                                     \
    if((LEN) == 1)                                                      \
      {                                                                 \
        NM_DATA_OP_APPLY_KIND_LEN(I, OP_KIND, START, COUNT, PTR, 1, P_OP); \
      }                                                                 \
    else if((LEN) == 2)                                                 \
      {                                                                 \
        NM_DATA_OP_APPLY_KIND_LEN(I, OP_KIND, START, COUNT, PTR, 2, P_OP); \
      }                                                                 \
    else if((LEN) == 4)                                                 \
      {                                                                 \
        NM_DATA_OP_APPLY_KIND_LEN(I, OP_KIND, START, COUNT, PTR, 4, P_OP); \
      }                                                                 \
    else if((LEN) == 8)                                                 \
      {                                                                 \
        NM_DATA_OP_APPLY_KIND_LEN(I, OP_KIND, START, COUNT, PTR, 8, P_OP); \
      }                                                                 \
    else if((LEN) == 16)                                                \
      {                                                                 \
        NM_DATA_OP_APPLY_KIND_LEN(I, OP_KIND, START, COUNT, PTR, 16, P_OP); \
      }                                                                 \
    else                                                                \
      {                                                                 \
        NM_DATA_OP_APPLY_KIND_LEN(I, OP_KIND, START, COUNT, PTR, (LEN), P_OP); \
      }                                                                 \
  }
 
#define NM_DATA_OP_APPLY_LOOP_OPT(I, COUNT, PTR, LEN, P_OP)             \
  {                                                                     \
    switch((P_OP)->kind)                                                \
      {                                                                 \
      case NM_DATA_OP_NOP:                                              \
        NM_DATA_OP_APPLY_KIND_LEN(I, nop, 0, (COUNT), (PTR), (LEN), (P_OP)); \
        break;                                                          \
      case NM_DATA_OP_COPY_FROM:                                        \
        NM_DATA_OP_APPLY_KIND_VECT(I, copy_from, 0, (COUNT), (PTR), (LEN), (P_OP)); \
        break;                                                          \
      case NM_DATA_OP_COPY_TO:                                          \
        NM_DATA_OP_APPLY_KIND_VECT(I, copy_to, 0, (COUNT), (PTR), (LEN), (P_OP)); \
        break;                                                          \
      case NM_DATA_OP_GETPROPS:                                         \
        NM_DATA_OP_APPLY_KIND_LEN(I, getprops, 0, (COUNT), (PTR), (LEN), (P_OP)); \
        break;                                                          \
      case NM_DATA_OP_DYNAMIC:                                          \
        NM_DATA_OP_APPLY_KIND_LEN(I, dynamic, 0, (COUNT), (PTR), (LEN), (P_OP)); \
        break;                                                          \
      case NM_DATA_OP_CHUNK:                                            \
        {                                                               \
          struct nm_data_op_context_chunk_s*__restrict__ p_chunk = &(P_OP)->context.chunk; \
          nm_len_t blocks_done = 0;                                     \
          if(p_chunk->done < p_chunk->chunk_offset)                     \
            {                                                           \
              nm_len_t skip_blocks = (p_chunk->chunk_offset - p_chunk->done) / (LEN); \
              if(skip_blocks > (COUNT)) skip_blocks = (COUNT);          \
              p_chunk->done += skip_blocks * (LEN);                     \
              blocks_done += skip_blocks;                               \
              if( (p_chunk->done < p_chunk->chunk_offset) &&            \
                  (blocks_done < (COUNT)) )                             \
                {                                                       \
                  assert(p_chunk->chunk_offset - p_chunk->done < (LEN)); \
                  const nm_len_t I = blocks_done;                       \
                  nm_data_op_apply_chunk((PTR), (LEN), p_chunk);        \
                  blocks_done++;                                        \
                  p_chunk->done += (LEN);                               \
                }                                                       \
            }                                                           \
          const nm_len_t chunk_end = p_chunk->chunk_offset + p_chunk->chunk_len; \
          if( (p_chunk->done < chunk_end) &&                            \
              (blocks_done < (COUNT)) )                                 \
            {                                                           \
              nm_len_t op_blocks = (chunk_end - p_chunk->done) / (LEN); \
              if(blocks_done + op_blocks > (COUNT))                     \
                op_blocks = (COUNT) - blocks_done;                      \
              switch(p_chunk->p_op->kind)                               \
                {                                                       \
                case NM_DATA_OP_NOP:                                    \
                  break;                                                \
                case NM_DATA_OP_COPY_FROM:                              \
                  NM_DATA_OP_APPLY_KIND_VECT(I, copy_from, blocks_done, op_blocks, (PTR), (LEN), p_chunk->p_op); \
                  break;                                                \
                case NM_DATA_OP_COPY_TO:                                \
                  NM_DATA_OP_APPLY_KIND_VECT(I, copy_to, blocks_done, op_blocks, (PTR), (LEN), p_chunk->p_op); \
                  break;                                                \
                case NM_DATA_OP_GETPROPS:                               \
                  NM_DATA_OP_APPLY_KIND_LEN(I, getprops, blocks_done, op_blocks, (PTR), (LEN), p_chunk->p_op); \
                  break;                                                \
                default:                                                \
                  {                                                     \
                    nm_len_t I;                                         \
                    for(I = blocks_done; I < blocks_done + op_blocks; I++) \
                      {                                                 \
                        nm_data_op_apply((PTR), (LEN), p_chunk->p_op);  \
                      }                                                 \
                  }                                                     \
                  break;                                                \
                }                                                       \
              blocks_done += op_blocks;                                 \
              p_chunk->done += op_blocks * (LEN);                       \
              if((p_chunk->done < chunk_end) && (blocks_done < (COUNT))) \
                {                                                       \
                  assert(chunk_end - p_chunk->done < (LEN));            \
                  const nm_len_t I = blocks_done;                       \
                  nm_data_op_apply_chunk((PTR), (LEN), p_chunk);        \
                  blocks_done++;                                        \
                  p_chunk->done += (LEN);                               \
                }                                                       \
            }                                                           \
          assert(blocks_done <= (COUNT));                               \
        }                                                               \
        break;                                                          \
      case NM_DATA_OP_SLICER_COROUTINE:                                 \
        {                                                               \
          struct nm_data_slicer_s*__restrict__ p_slicer = (P_OP)->context.slicer_coroutine.p_slicer; \
          nm_len_t blocks_done = 0;                                     \
          while(blocks_done < (COUNT))                                  \
            {                                                           \
                const nm_len_t slice_len = p_slicer->coroutine.slice_len; \
                nm_len_t slice_blocks = slice_len / (LEN);              \
                if(blocks_done + slice_blocks > (COUNT))                \
                  slice_blocks = (COUNT) - blocks_done;                 \
                if(slice_blocks > 0)                                    \
                  {                                                     \
                    switch(p_slicer->coroutine.op.kind)                 \
                      {                                                 \
                      case NM_DATA_OP_NOP:                              \
                        NM_DATA_OP_APPLY_KIND_LEN(I, nop, blocks_done, slice_blocks, (PTR), (LEN), &p_slicer->coroutine.op); \
                        break;                                          \
                      case NM_DATA_OP_COPY_FROM:                        \
                        NM_DATA_OP_APPLY_KIND_VECT(I, copy_from, blocks_done, slice_blocks, (PTR), (LEN), &p_slicer->coroutine.op); \
                        break;                                          \
                      case NM_DATA_OP_COPY_TO:                          \
                        NM_DATA_OP_APPLY_KIND_VECT(I, copy_to, blocks_done, slice_blocks, (PTR), (LEN), &p_slicer->coroutine.op); \
                        break;                                          \
                      default:                                          \
                        {                                               \
                          nm_len_t I;                                   \
                          for(I = blocks_done; I < blocks_done + slice_blocks; I++) \
                            {                                           \
                              nm_data_op_apply((PTR), (LEN), &p_slicer->coroutine.op); \
                            }                                           \
                        }                                               \
                        break;                                          \
                      }                                                 \
                    p_slicer->coroutine.slice_len = slice_len - slice_blocks * (LEN); \
                    blocks_done += slice_blocks;                        \
                    if(slice_len == slice_blocks * (LEN))               \
                      {                                                 \
                        nm_data_coroutine_yield_to_caller(&p_slicer->coroutine.coroutine); \
                      }                                                 \
                    else if(blocks_done < (COUNT))                      \
                      {                                                 \
                        const nm_len_t I = blocks_done;                 \
                        nm_data_op_apply_slicer_coroutine((PTR), (LEN), &(P_OP)->context.slicer_coroutine); \
                        blocks_done++;                                  \
                      }                                                 \
                  }                                                     \
                else if(blocks_done < (COUNT))                          \
                  {                                                     \
                    const nm_len_t I = blocks_done;                     \
                    nm_data_op_apply_slicer_coroutine((PTR), (LEN), &(P_OP)->context.slicer_coroutine); \
                    blocks_done++;                                      \
                  }                                                     \
              }                                                         \
          assert(blocks_done <= (COUNT));                               \
        }                                                               \
        break;                                                          \
      default:                                                          \
        {                                                               \
          nm_len_t I;                                                   \
          for(I = 0; I < (COUNT); I++)                                  \
            {                                                           \
              nm_data_op_apply((PTR), (LEN), (P_OP));                   \
            }                                                           \
        }                                                               \
        break;                                                          \
      }                                                                 \
  }
 
#define NM_DATA_OP_APPLY_LOOP NM_DATA_OP_APPLY_LOOP_OPT
 
 
/* ** Built-in data types ********************************** */
 
struct nm_data_null_s
{
  int dummy; 
};
extern const struct nm_data_ops_s nm_data_ops_null;
NM_DATA_TYPE(null, struct nm_data_null_s, &nm_data_ops_null);
static inline void nm_data_null_build(struct nm_data_s*p_data)
{
  struct nm_data_null_s n = { 0 };
  nm_data_null_set(p_data, n);
}
static inline int nm_data_isnull(struct nm_data_s*p_data)
{
  return (p_data->ops.p_traversal == nm_data_ops_null.p_traversal);
}
 
struct nm_data_contiguous_s
{
  void*ptr;     
  nm_len_t len; 
};
extern const struct nm_data_ops_s nm_data_ops_contiguous;
NM_DATA_TYPE(contiguous, struct nm_data_contiguous_s, &nm_data_ops_contiguous);
 
static inline void nm_data_contiguous_build(struct nm_data_s*p_data, void*ptr, nm_len_t len)
{
  struct nm_data_contiguous_s dc;
  dc.ptr = ptr;
  dc.len = len;
  nm_data_contiguous_set(p_data, dc);
}
 
struct nm_data_iov_s
{
  const struct iovec*v;
  int n;
};
extern const struct nm_data_ops_s nm_data_ops_iov;
NM_DATA_TYPE(iov, struct nm_data_iov_s, &nm_data_ops_iov);
 
static inline void nm_data_iov_build(struct nm_data_s*p_data, const struct iovec*v, int n)
{
  struct nm_data_iov_s di;
  di.v = v;
  di.n = n;
  nm_data_iov_set(p_data, di);
}
 
struct nm_data_datav_s
{
  struct nm_datav_s*p_datav;
};
extern const struct nm_data_ops_s nm_data_ops_datav;
NM_DATA_TYPE(datav, struct nm_data_datav_s, &nm_data_ops_datav);
 
static inline void nm_data_datav_build(struct nm_data_s*p_datav_data, struct nm_datav_s*p_datav)
{
  p_datav->commited = 1;
  struct nm_data_datav_s dv;
  dv.p_datav = p_datav;
  nm_data_datav_set(p_datav_data, dv);
}
 
struct nm_data_excerpt_s
{
  nm_len_t chunk_offset;
  nm_len_t chunk_len;
  struct nm_data_s*p_data;
};
extern const struct nm_data_ops_s nm_data_ops_excerpt;
NM_DATA_TYPE(excerpt, struct nm_data_excerpt_s, &nm_data_ops_excerpt);
 
static inline void nm_data_excerpt_build(struct nm_data_s*p_data, struct nm_data_s*p_inner_data,
                                         nm_len_t chunk_offset, nm_len_t chunk_len)
{
  struct nm_data_excerpt_s de;
  de.chunk_offset = chunk_offset;
  de.chunk_len = chunk_len;
  de.p_data = p_inner_data;
  nm_data_excerpt_set(p_data, de);
}
 
 
/* ** Helper functions ************************************* */
 
static inline void nm_data_traversal_op_apply(const struct nm_data_s*p_data, struct nm_data_op_s*p_op)
{
  assert(p_data->ops.p_traversal != NULL);
  (*p_data->ops.p_traversal)((void*)p_data->_content, p_op);
}
 
static inline void nm_data_traversal_apply(const struct nm_data_s*p_data, nm_data_apply_t p_apply, void*_context)
{
  struct nm_data_op_s op;
  op.kind = NM_DATA_OP_DYNAMIC;
  op.context.dynamic.p_apply = p_apply;
  op.context.dynamic.p_apply_context = _context;
  nm_data_traversal_op_apply(p_data, &op);
}
 
void nm_data_chunk_extractor_op_traversal(const struct nm_data_s*p_data, nm_len_t chunk_offset, nm_len_t chunk_len,
                                          struct nm_data_op_s*p_op);
 
void nm_data_chunk_extractor_traversal(const struct nm_data_s*p_data, nm_len_t chunk_offset, nm_len_t chunk_len,
                                       nm_data_apply_t p_apply, void*p_apply_context);
 
static inline const struct nm_data_properties_s*nm_data_properties_get(const struct nm_data_s*p_data)
{
  return &p_data->props;
}
 
static inline nm_len_t nm_data_size(const struct nm_data_s*p_data)
{
  const struct nm_data_properties_s*p_props = nm_data_properties_get((struct nm_data_s*)p_data);
  return p_props->size;
}
 
void*nm_data_baseptr_get(const struct nm_data_s*p_data);
 
void*nm_data_chunk_baseptr_get(const struct nm_data_s*p_data, nm_len_t chunk_offset, nm_len_t chunk_len);
 
void nm_data_chunk_properties_compute(const struct nm_data_s*p_data, nm_len_t chunk_offset, nm_len_t chunk_len,
                                      struct nm_data_properties_s*p_props);
 
nm_len_t nm_data_chunk_first_get(const struct nm_data_s*p_data, nm_len_t chunk_offset, nm_len_t chunk_len, int n);
 
uint32_t nm_data_checksum(const struct nm_data_s*p_data);
 
void nm_data_copy_from(const struct nm_data_s*p_data, nm_len_t offset, nm_len_t len, void*destbuf);
 
void nm_data_copy_to(const struct nm_data_s*p_data, nm_len_t offset, nm_len_t len, const void*srcbuf);
 
void nm_data_copy(struct nm_data_s*p_dest, struct nm_data_s*p_from);
 
/* ** Data slicer ****************************************** */
 
/* ** coroutine */
 
struct nm_data_coroutine_s;
 
typedef void (*nm_data_coroutine_worker_t)(struct nm_data_coroutine_s*p_coroutine, void*_user_data);
 
enum nm_data_coroutine_kind_e
  {
    NM_DATA_COROUTINE_NONE,
    NM_DATA_COROUTINE_UCONTEXT, 
    NM_DATA_COROUTINE_LONGJMP   
  };
typedef enum nm_data_coroutine_kind_e nm_data_coroutine_kind_t;
 
struct nm_data_coroutine_s
{
  nm_data_coroutine_kind_t kind;
  union
  {
    struct nm_data_coroutine_ucontext_s*p_ucontext;
    struct nm_data_coroutine_longjmp_s*p_longjmp;
  } impl;
  nm_data_coroutine_worker_t p_worker;
  void*p_worker_data;
};
 
void nm_data_coroutine_yield_to_caller(struct nm_data_coroutine_s*p_coroutine);
void nm_data_coroutine_yield_to_data(struct nm_data_coroutine_s*p_coroutine);
 
/* ** slicer */
 
typedef enum
  {
    NM_DATA_SLICER_NONE = 0,
    NM_DATA_SLICER_COROUTINE, 
    NM_DATA_SLICER_CONTIG     
  } nm_data_slicer_kind_t;
 
typedef struct nm_data_slicer_s
{
  nm_data_slicer_kind_t kind;
  const struct nm_data_s*p_data;
  nm_len_t done;                    
  union
  {
    struct
    {
      struct nm_data_op_s op;       
      volatile nm_len_t slice_len;  
      struct nm_data_coroutine_s coroutine;
    } coroutine;                    
    struct
    {
      void*p_baseptr;               
    } contig;                       
  };
} nm_data_slicer_t;
 
#define NM_DATA_SLICER_NULL ((struct nm_data_slicer_s){ .kind = NM_DATA_SLICER_NONE })
 
static inline int nm_data_slicer_isnull(const nm_data_slicer_t*p_slicer)
{
  return (p_slicer->kind == NM_DATA_SLICER_NONE);
}
 
void nm_data_slicer_init(nm_data_slicer_t*p_slicer, const struct nm_data_s*p_data);
void nm_data_slicer_copy_from(nm_data_slicer_t*p_slicer, void*dest_ptr, nm_len_t slice_len);
void nm_data_slicer_copy_to(nm_data_slicer_t*p_slicer, const void*src_ptr, nm_len_t slice_len);
void nm_data_slicer_forward(nm_data_slicer_t*p_slicer, nm_len_t offset);
void nm_data_slicer_destroy(nm_data_slicer_t*p_slicer);
 
/* ********************************************************* */
/* ** inline functions */
 
 
/* ** datav */
 
static inline void nm_datav_init(struct nm_datav_s*p_datav)
{
  p_datav->p_data = &p_datav->data[0];
  p_datav->n_data = 0;
  p_datav->allocated = 0;
  p_datav->commited = 0;
}
 
static inline void nm_datav_destroy(struct nm_datav_s*p_datav)
{
  assert(p_datav->p_data != NULL);
  if(p_datav->p_data != &p_datav->data[0])
   {
     padico_free(p_datav->p_data);
     p_datav->p_data = NULL;
   }
 }
 
static inline void nm_datav_add_chunk_data(struct nm_datav_s*p_datav, const struct nm_data_s*p_data)
{
  assert(!p_datav->commited); /* cannot modify datav once it is used as a nm_data */
  if(p_datav->n_data == NM_DATAV_INIT_SIZE)
    {
      assert(p_datav->p_data == &p_datav->data[0]);
      p_datav->allocated = NM_DATAV_INIT_SIZE * 2;
      p_datav->p_data = (struct nm_data_s*)padico_malloc(p_datav->allocated * sizeof(struct nm_data_s));
      memcpy(p_datav->p_data, &p_datav->data[0], p_datav->n_data * sizeof(struct nm_data_s));
    }
  else if((p_datav->n_data > NM_DATAV_INIT_SIZE) &&
          (p_datav->n_data > p_datav->allocated - 1))
    {
      assert(p_datav->p_data != &p_datav->data[0]);
      p_datav->allocated *= 2;
      p_datav->p_data = (struct nm_data_s*)padico_realloc(p_datav->p_data, p_datav->allocated * sizeof(struct nm_data_s));
    }
  p_datav->p_data[p_datav->n_data] = *p_data;
  p_datav->n_data++;
}
 
static inline void nm_datav_add_chunk(struct nm_datav_s*p_datav, const void*ptr, nm_len_t len)
{
  struct nm_data_s data;
  nm_data_contiguous_build(&data, (void*)ptr, len);
  nm_datav_add_chunk_data(p_datav, &data);
}
 
static inline nm_len_t nm_datav_size(struct nm_datav_s*p_datav)
{
  nm_len_t size = 0;
  int i;
  for(i = 0; i < p_datav->n_data; i++)
    {
      size += nm_data_size(&p_datav->p_data[i]);
    }
  return size;
}
 
static inline void nm_datav_uncommit(struct nm_datav_s*p_datav)
{
  assert(p_datav->commited);
  p_datav->commited = 0;
}
 
/* ** GPU */
 
static inline void nm_data_propertie_gpu_preinit(struct nm_data_properties_s*p_props __attribute__((unused)))
{
#ifdef NMAD_CUDA
  p_props->is_cuda = -1;
#endif /* NMAD_CUDA */
#ifdef NMAD_HIP
  p_props->is_hip = -1;
#endif /* NMAD_HIP */
}
 
static inline void nm_data_propertie_gpu_postinit(const struct nm_data_properties_s*p_props __attribute__((unused)))
{
#ifdef NMAD_CUDA
  if(p_props->is_cuda == -1)
    {
      NM_FATAL("incorrect data properties: is_cuda = %d (should be either 0 or 1)\n", p_props->is_cuda);
    }
#endif /* NMAD_CUDA */
#ifdef NMAD_HIP
  if(p_props->is_hip == -1)
    {
      NM_FATAL("incorrect data properties: is_hip = %d (should be either 0 or 1)\n", p_props->is_hip);
    }
#endif /* NMAD_HIP */
}
 
 
static inline void nm_data_memcpy_from(void*p_dest, const void*p_src, nm_len_t len, const struct nm_data_properties_s*p_props)
{
#if defined(NMAD_CUDA)
  if(p_props->is_cuda)
    {
      cudaError_t rc = cudaMemcpy(p_dest /* host */, p_src /* gpu */, len, cudaMemcpyDeviceToHost);
      if(rc != cudaSuccess)
        {
          NM_FATAL("CUDA cannot copy %ld bytes from GPU; error %s\n", len, cudaGetErrorString(rc));
        }
      return;
    }
#endif /* NMAD_CUDA */
 
#if defined(NMAD_HIP)
  if(p_props->is_hip)
    {
      hipError_t rc = hipMemcpy(p_dest /* host */, p_src /* gpu */, len, hipMemcpyDeviceToHost);
      if(rc != hipSuccess)
        {
          NM_FATAL("HIP cannot copy %ld bytes from GPU; error %s\n", len, hipGetErrorString(rc));
        }
      return;
    }
#endif /* NMAD_HIP */
 
  /* default case: host memory */
  memcpy(p_dest, p_src, len);
}
 
static inline void nm_data_memcpy_to(void*p_dest, const void*p_src, nm_len_t len, const struct nm_data_properties_s*p_props)
{
#if defined(NMAD_CUDA)
  if(p_props->is_cuda)
    {
      cudaError_t rc = cudaMemcpy(p_dest /* gpu */, p_src /* src */, len, cudaMemcpyHostToDevice);
      if(rc != cudaSuccess)
        {
          NM_FATAL("CUDA cannot copy %ld bytes from GPU; error %s\n", len, cudaGetErrorString(rc));
        }
      return;
    }
#endif /* NMAD_CUDA */
 
#if defined(NMAD_HIP)
  if(p_props->is_hip)
    {
      hipError_t rc = hipMemcpy(p_dest /* gpu */, p_src /* host */, len, hipMemcpyHostToDevice);
      if(rc != hipSuccess)
        {
          NM_FATAL("HIP cannot copy %ld bytes from GPU; error %s\n", len, hipGetErrorString(rc));
        }
      return;
    }
#endif /* NMAD_HIP */
 
  /* default case: host memory */
  memcpy(p_dest, p_src, len);
}
 
 
/* ** data op inline */
 
static inline void nm_data_op_apply_nop(void*__restrict__ p_ptr, nm_len_t len, struct nm_data_op_context_nop_s*__restrict__ p_nop)
{
  /* empty */
}
 
static inline void nm_data_op_apply_copy_from(void*__restrict__ p_ptr, nm_len_t len, struct nm_data_op_context_copy_from_s*__restrict__ p_copy_from)
{
  nm_data_memcpy_from(p_copy_from->p_dest_ptr, p_ptr, len, &p_copy_from->p_data->props);
  p_copy_from->p_dest_ptr = (char*)p_copy_from->p_dest_ptr + len;
}
 
static inline void nm_data_op_apply_copy_to(void*__restrict__ p_ptr, nm_len_t len, struct nm_data_op_context_copy_to_s*__restrict__ p_copy_to)
{
  nm_data_memcpy_to(p_ptr, p_copy_to->p_src_ptr, len, &p_copy_to->p_data->props);
  p_copy_to->p_src_ptr = (char*)p_copy_to->p_src_ptr + len;
}
 
static inline void nm_data_op_apply_dynamic(void*__restrict__ p_ptr, nm_len_t len, struct nm_data_op_context_dynamic_s*__restrict__ p_dynamic)
{
  (*p_dynamic->p_apply)(p_ptr, len, p_dynamic->p_apply_context);
}
 
static inline void nm_data_op_apply_getprops(void*__restrict__ p_ptr, nm_len_t len, struct nm_data_op_context_getprops_s*__restrict__ p_getprops)
{
  p_getprops->props.size += len;
  p_getprops->props.blocks += 1;
  if(p_getprops->props.is_contig)
    {
      if((p_getprops->p_blockend != NULL) && (p_ptr != p_getprops->p_blockend))
        p_getprops->props.is_contig = 0;
      p_getprops->p_blockend = (char*)p_ptr + len;
#if defined(NMAD_CUDA) || defined(NMAD_HIP)
      if(p_getprops->p_blockend == NULL) /* test only first chunk */
        {
          nm_data_properties_gpu_fill(&p_getprops->props, p_ptr);
        }
#endif /* CUDA || HIP */
    }
}
 
static inline void nm_data_op_apply_chunk(void*__restrict__ p_ptr, nm_len_t len, struct nm_data_op_context_chunk_s*__restrict__ p_chunk)
{
  const nm_len_t chunk_offset = p_chunk->chunk_offset;
  const nm_len_t chunk_len = p_chunk->chunk_len;
  if( (!(p_chunk->done + len <= p_chunk->chunk_offset))                    /* data before chunk- do nothing */
      &&
      (!(p_chunk->done >= p_chunk->chunk_offset + p_chunk->chunk_len))) /* data after chunk- do nothing */
    {
      /* data in chunk */
      const nm_len_t block_offset = (p_chunk->done < chunk_offset) ? (chunk_offset - p_chunk->done) : 0;
      const nm_len_t block_len = (chunk_offset + chunk_len > p_chunk->done + len) ?
        (len - block_offset) : (chunk_offset + chunk_len - p_chunk->done - block_offset);
      nm_data_op_apply((char*)p_ptr + block_offset, block_len, p_chunk->p_op);
    }
  p_chunk->done += len;
}
 
static inline void nm_data_op_apply_slicer_coroutine(void*__restrict__ p_ptr, nm_len_t len, struct nm_data_op_context_slicer_coroutine_s*__restrict__ p_slicer_coroutine)
{
  struct nm_data_slicer_s*__restrict__ p_slicer = p_slicer_coroutine->p_slicer;
 restart:
  ;
  const nm_len_t slice_len = p_slicer->coroutine.slice_len; /* take a snapshot of volatile var */
  const nm_len_t chunk_len = (len > slice_len) ? slice_len : len;
  nm_data_op_apply(p_ptr, chunk_len, &p_slicer->coroutine.op);
  if(slice_len == chunk_len)
    {
      /* slice is done- give hand back to caller context */
      nm_data_coroutine_yield_to_caller(&p_slicer->coroutine.coroutine);
      /* back from caller context */
      if(len != chunk_len)
        {
          /* crop block and process remainder */
          p_ptr = (char*)p_ptr + chunk_len;
          len -= chunk_len;
          goto restart;
        }
    }
  else
    {
      p_slicer->coroutine.slice_len = slice_len - chunk_len;
    }
}
 
static inline void nm_data_op_apply(void*__restrict__ p_ptr, nm_len_t len, struct nm_data_op_s*__restrict__ p_op)
{
  switch(p_op->kind)
    {
    case NM_DATA_OP_NOP:
      nm_data_op_apply_nop(p_ptr, len, &p_op->context.nop);
      break;
    case NM_DATA_OP_COPY_FROM:
      nm_data_op_apply_copy_from(p_ptr, len, &p_op->context.copy_from);
      break;
    case NM_DATA_OP_COPY_TO:
      nm_data_op_apply_copy_to(p_ptr, len, &p_op->context.copy_to);
      break;
    case NM_DATA_OP_DYNAMIC:
      nm_data_op_apply_dynamic(p_ptr, len, &p_op->context.dynamic);
      break;
    case NM_DATA_OP_GETPROPS:
      nm_data_op_apply_getprops(p_ptr, len, &p_op->context.getprops);
      break;
    case NM_DATA_OP_CHUNK:
      nm_data_op_apply_chunk(p_ptr, len, &p_op->context.chunk);
      break;
    case NM_DATA_OP_SLICER_COROUTINE:
      nm_data_op_apply_slicer_coroutine(p_ptr, len, &p_op->context.slicer_coroutine);
      break;
    case NM_DATA_OP_NONE:
      NM_FATAL("cannot apply NM_DATA_OP_NONE.\n");
      break;
    default:
      NM_FATAL("op = %d not managed in apply.\n", p_op->kind);
      break;
    }
}
 
 
#endif /* NM_DATA_H */