#ifndef _MY_FRAME_TYPE_H_
#define _MY_FRAME_TYPE_H_

#ifdef _CMEM_USE_GPU_
    #if (defined(_USE_CUDA_) || defined(_USE_ROCM_))
        #define FUNTYPE __global__
        #define FUNTYPE_H __global__     
           
        #define _GMEM_   
        #define _DEF_SMEM_ __shared__
        #define _LMEM_      
    #else
        #define FUNTYPE __kernel
        #define FUNTYPE_H       
           
        #define _GMEM_ __global         
        #define _DEF_SMEM_ __local
        #define _LMEM_
        
        #define __device__
        #define __host__      
        
    #endif
#else
    #define FUNTYPE
    #define FUNTYPE_H

    #define __host__
    #define __device__
    
    #define _GMEM_
    #define _LMEM_
    #define _DEF_SMEM_
#endif

#ifdef _USE_OPENCL_  
   #define _DSMEM_ _DEF_SMEM_  
   #define _SMEM_ _DEF_SMEM_   
#else   
    #define _SMEM_   
    #define _DSMEM_
#endif
 
////////////////////////////////////

// FP8_16 格式定义
#define FP8_E4M3_NAN      0x7F     // NaN (all exponent bits set, mantissa non-zero)
#define FP8_E4M3_INF      0x7E     // Infinity (max exponent, zero mantissa)
#define FP8_E5M2_NAN      0x7F     // NaN
#define FP8_E5M2_INF      0x7C     // Infinity
#define FP8_E5M2_MAX  0b01111011   // 最大有限值 (57344)
#define FP32_NAN 0x7FC00000
#define FP32_INF 0x7F800000

#define MyVoid float
//#define MyVoid void

typedef unsigned char fp8_e4m3;  // 1 sign bit (optional), 4 exponent bits, 3 mantissa bits
typedef unsigned char fp8_e5m2;  // 1 sign bit (optional), 5 exponent bits, 2 mantissa bits
typedef unsigned short fp16;
typedef float fp32; 

typedef struct {
    float x, y, z, w;
} CCNNfloat4;

typedef struct {
    short x, y, z, w;
} CCNNshort4;

union FP32_vector4 {
   CCNNfloat4 f;
   fp32 unit[4];
};

union FP8_E4M3_vector4 {
   float f;
   fp8_e4m3 unit[4];
};

union FP8_E5M2_vector4 {
   float f;
   fp8_e5m2 unit[4];
};

union FP16_vector4 {
   CCNNshort4 f;
   fp16 unit[4];
};

//Active function type
enum ActiveType {SIGMOD, RELU, LRELU, TANH};

enum LayerType {INPT, CONV, DECONV, BIAS, AVEPOOL, AVEUS, MAXPOOL, MAXUS, ACTIVE, LN, BN, L2N, FC, \
                CROP, EXPAND, SFMIMG, SFMVEC, SSD, REGRESS, DROPOUT, GRU, CBA, FBA, RESU, RESN, MFL,\
                BTNK1, BTNK2, YOLO3, Linear, F2L, SFM, ATTEN, SUM, SEQ2SEQ, RNN, C_SCALE, SEQ_SUM_CAT,\
                T2L, FILTER, MULTI_HEAD_ATTEN, TRANSFORMER_ENCODE, TRANSFORMER_DECODE, TRANSFORMER};

//CPU Multi-thread Action Type
enum ThreadComType {COMPUTERY, COMPUTERD, UPDATEK};

enum MemType {GPU_TYPE, CPU_TYPE};

//Conv type     0     1     2
enum ConvType{FULL, SAME, VALID};

//Softmax lost Type
enum SLTYPE {SF, FL};

// in Data to next layer method
enum DATA_FUSION_TYPE {DATA_ADD, DATA_DOT, DATA_CAT_N_HW_C, DATA_CAT_N_C_HW, DATA_NC_CHW};

enum DATA_SUM_TYPE {CHW_SUM, C_SUM, HW_SUM};

enum FLOAT_TYPE {FP32_TYPE, FP8_E4M3_TYPE, FP8_E5M2_TYPE, FP16_TYPE};

#define OPENCL_KERNEL_PATH_OPTION "-I ./inc"
#define CONST_STR_LEN 1024

#define MY_THREAD_PER_BLOCK 256
#define NATURAL_NUM 2.71828F
#define EPISION 0.00001F
#define PI 3.14159F

#define CONST_ONE_DIM_SPLIT_TWO_DIM_X_DIM 64
#define CONST_CHANNEL_NUM 4
#define THREADS_PER_BLOCK_FOR_REDUCTION 256
#define TILE 16
#define MY_LOOP(i, basic, N) for (i = basic; i < basic + N; i++)
#define SCH_4_DIM 4
#define SCH_3_DIM 3
#define SCH_2_DIM 2
#define FC_LOST_GAMMA 2
#define FC_LOST_ALPHA 0.25F
#define NEGATIVE_MAX_NUM -100000000.0F

#define SSD_THREADS_NUM 256
#define SSD_BLOCKS_NUM 1
#define SECOND_RULE_IOU 0.5F

#ifdef _CMEM_USE_CPU_
    #define SCH_DIM SCH_2_DIM
#else
    #define SCH_DIM SCH_4_DIM
#endif

#define MapReadBegin(fType, devID, p, l, xp) MyVoid *xp = XPUMapReadBegin(fType, devID, p, l)
#define MapWriteBegin(fType, devID, p, l, xp) MyVoid *xp = XPUMapWriteBegin(fType, devID, p, l)
#define MapReadWriteBegin(fType, devID, p, l, xp) MyVoid *xp = XPUMapReadWriteBegin(fType, devID, p, l)
#define MapReadEnd(fType, devID, p, l, xp) XPUMapReadEnd(fType, devID, p, l, xp)
#define MapWriteEnd(fType, devID, p, l, xp) XPUMapWriteEnd(fType, devID, p, l, xp)
#define MapReadWriteEnd(fType, devID, p, l, xp) XPUMapReadWriteEnd(fType, devID, p, l, xp)
 
#ifdef _USE_OPENCL_
    #define RegisterDeclareVoidPoint(t, para) char share##para[4];MyVoid *para=(MyVoid *)share##para; 
#else   
    //#define RegisterDeclareVoidPoint(t, para) MyVoid *para; if (t==0) {char share##para[4];para=(MyVoid *)share##para;} else if (t<=2) {char share##para[1];para=(MyVoid *)share##para;} else {char share##para[2];para=(MyVoid *)share##para;}    
    #define RegisterDeclareVoidPoint(t, para) float share##para;MyVoid *para=(MyVoid *)&share##para; 
#endif

#define RegisterDeclareVoidPointInitFloat(t, para, fNum) RegisterDeclareVoidPoint(t, para);SetMyVoidValueByFloat(t, para, fNum); 
#define RegisterDeclareVoidPointStackArray(t, para, size) float share##para[size];MyVoid *para=(MyVoid *)share##para;

#define B_M 128
#define B_N 128
#define B_K 8
#define K_G 1 

#endif
