#ifndef MATFP16_KERNEL_BODY_H_
#define MATFP16_KERNEL_BODY_H_

__host__ __device__ fp16 GetSoftMaxGrad_FP16(int lostType, int cIndex, int trueChannelIndex, fp16 P_t, fp16 P_i, float coff);
__host__ __device__ fp16 GetSoftMaxLoss_FP16(int lostType, fp16 P_t, float coff);
__host__ __device__ int VectorMaxIndex_FP16(fp16 *arrayDataSrc, int veclength);
__host__ __device__ void LogSumExp_FP16(fp16 *result, fp16 *arrFloat, int len, int tIndex);
__host__ __device__ void GetSoftMaxLossBySrcOut_FP16(fp16 *result, int lostType, fp16 *arrFloat, int len, int tIndex, fp16 coff);

#ifdef _CMEM_USE_GPU_
__device__ void AtomicFloat4AddIn_FP16(fp16 *dst, fp16 *src);
#endif

__host__ __device__ union FP16_vector4 FP16_FLOAT4_Multi(union FP16_vector4 f1, fp16 coff);
__host__ __device__ union FP16_vector4 FP16_FLOAT4_ADD_Multi(union FP16_vector4 f, union FP16_vector4 f1, fp16 coff);
   
__host__ __device__ fp16 ActiveFun_FP16(fp16 x, int aType);

__host__ __device__ fp16 DeActiveFun_FP16(fp16 y, int aType);

__host__ __device__ void MaxUpSampleYKernelFun_FP16(_GMEM_ fp16 *Y, _GMEM_ fp16 *X, _GMEM_ fp16 *rMat, _GMEM_ fp16 *cMat, \
                                               int kernelSize, int C, int H, int W, int nIdx, int cIdx, int hIdx, int wIdx, int rtPad, int rbPad, \
                                               int clPad, int crPad);

__host__ __device__ void MaxPoolYKernelFun_FP16(_GMEM_ fp16 *Y, _GMEM_ fp16 *X, _GMEM_ fp16 *rMat, _GMEM_ fp16 *cMat, \
                                           int kernelSize, int C, int H, int W, int nIdx, int cIdx, int hIdx, int wIdx, int rtPad, int rbPad, \
                                           int clPad, int crPad);

__host__ __device__ void MaxUpSampleDKernelFun_FP16(_GMEM_ fp16 *rD, _GMEM_ fp16 *D, _GMEM_ fp16 *rMat, _GMEM_ fp16 *cMat, \
                                               int kernelSize, int C, int H, int W, int nIdx, int cIdx, int hIdx, int wIdx, int rtPad, int rbPad, \
                                               int clPad, int crPad);

__host__ __device__ void MaxPoolDKernelFun_FP16(_GMEM_ fp16 *rD, _GMEM_ fp16 *D, _GMEM_ fp16 *rMat, _GMEM_ fp16 *cMat, \
                                           int kernelSize, int C, int H, int W, int nIdx, int cIdx, int hIdx, int wIdx, int rtPad, int rbPad, \
                                           int clPad, int crPad);

__host__ __device__ void MatrixAveupSample_FP16(_GMEM_ fp16 *Y, _GMEM_ fp16 *X, int kernelSize, int C, int H, int W, int nIdx, \
                                           int cIdx, int hIdx, int wIdx, int rtPad, int rbPad, int clPad, int crPad, fp16 f);

__host__ __device__ void MatrixAvePooling_FP16(_GMEM_ fp16 *Y, _GMEM_ fp16 *X, int kernelSize, int C, int H, int W, int nIdx, int cIdx, \
                                          int hIdx, int wIdx, int rtPad, int rbPad, int clPad, int crPad, fp16 f);

#ifdef _CMEM_USE_GPU_
#if (defined(_USE_CUDA_) || defined(_USE_ROCM_))
FUNTYPE void ArrayAddArrayKernel_FP16(fp16 *C, fp16 *A, fp16 *B, int len);
#endif
#endif

#ifdef _CMEM_USE_GPU_
#if (defined(_USE_CUDA_) || defined(_USE_ROCM_))
FUNTYPE void ArraySubArrayKernel_FP16(fp16 *C, fp16 *A, fp16 *B, int len);
#endif
#endif

#ifdef _CMEM_USE_GPU_
#if (defined(_USE_CUDA_) || defined(_USE_ROCM_))
FUNTYPE void ArrayMultiFactorAddBiasGPUKernel_FP16(fp16 *res, fp16 *src, fp16 *factor, fp16 *bias, int len);
#endif
#endif

#ifdef _CMEM_USE_GPU_
#if (defined(_USE_CUDA_) || defined(_USE_ROCM_))
FUNTYPE void ArrayMultiFactorGPUKernel_FP16(fp16 *res, fp16 *src, fp16 *factor, int len);
#endif
#endif

#ifdef _CMEM_USE_GPU_
#if (defined(_USE_CUDA_) || defined(_USE_ROCM_))
FUNTYPE void ArrayDotArrayKernelGPU_FP16(fp16 *resultArray, fp16 *inArray1, fp16 *inArray2, int len);
#endif
#endif

#ifdef _CMEM_USE_GPU_
#if (defined(_USE_CUDA_) || defined(_USE_ROCM_))
FUNTYPE void ArrayDivArrayKernelGPU_FP16(fp16 *resultArray, fp16 *inArray1, fp16 *inArray2, int len);
#endif
#endif

#ifdef _CMEM_USE_CPU_
void ActiveFunKernel_FP16(fp16 *R, fp16 *A, int dSize, int mFUNTYPE, int mainGlobalThreadId);
#else
#if (defined(_USE_CUDA_) || defined(_USE_ROCM_))
FUNTYPE void ActiveFunKernel_FP16(fp16 *R, fp16 *A, int dSize, int mFUNTYPE);
#endif
#endif

#ifdef _CMEM_USE_CPU_
void DeActiveFunKernel_FP16(fp16 *R, fp16 *D, fp16 *A, int dSize, int mFUNTYPE, int mainGlobalThreadId);
#else
#if (defined(_USE_CUDA_) || defined(_USE_ROCM_))
FUNTYPE void DeActiveFunKernel_FP16(fp16 *R, fp16 *D, fp16 *A, int dSize, int mFUNTYPE);
#endif
#endif

#ifdef _CMEM_USE_GPU_
#if (defined(_USE_CUDA_) || defined(_USE_ROCM_))
FUNTYPE void MatMultiKernelByTile_FP16(fp16 *A_d, fp16 *B_d, fp16 *C_dOff, \
                                  int offset, int m, int k, int n, int ifCAddin);
#endif
#endif

#ifdef _CMEM_USE_GPU_
__device__ void FloatBuffSet4ByteFloat_FP16(fp16 *dst, fp16 *src);
#endif

#ifdef _CMEM_USE_GPU_
__device__ void StoreRegCToC_FP16(int cH, int cW, int ldc, fp16 regC[][8], fp16 *C_d, int tmpBlockIdxX, int tmpBlockIdxY, int threadIdxX, int ifCAddin);
#endif

#ifdef _CMEM_USE_GPU_
__device__ void MatOuterMulti64_FP16(fp16 iregC[][8], fp16 *iregA, fp16 *iregB, int index);
#endif

#ifdef _CMEM_USE_GPU_
#if (defined(_USE_CUDA_) || defined(_USE_ROCM_))
__device__ void LoadABRegFromShareMem_FP16(int switchIdx, int smallIdx, fp16 regA[2][8], fp16 regB[2][8], \
                                      fp16 smA[2][8][128], fp16 smB[2][8][128], \
                                      int tmpThreadIdxX);
#endif
#endif

#ifdef _CMEM_USE_GPU_
#if (defined(_USE_CUDA_) || defined(_USE_ROCM_))
__device__ void SetFloat4Data_FP16(union FP16_vector4 *dst, fp16 *src);
#endif
#endif

#ifdef _CMEM_USE_GPU_
__device__ union FP16_vector4 GetFloat4Data_FP16(fp16 *src);
#endif

#ifdef _CMEM_USE_GPU_ 
#if (defined(_USE_CUDA_) || defined(_USE_ROCM_))
__device__ void Trans4FloatSaveToFloatBuff_FP16(fp16 *dst, int f4offset, int index, fp16 inData);
#endif
#endif

#ifdef _CMEM_USE_GPU_
#if (defined(_USE_CUDA_) || defined(_USE_ROCM_))
__device__ void LoadABShareMemFromDevMem_FP16(int switchIdx, fp16 smA[2][8][128], \
                                         fp16 smB[2][8][128], int m, int n, int k, int t, \
                                         int tmpBlockIdxX, int tmpBlockIdxY, int tmpThreadIdxX, \
                                         fp16 *A_df, fp16 *B_df);
#endif
#endif

#ifdef _CMEM_USE_GPU_
#if (defined(_USE_CUDA_) || defined(_USE_ROCM_))
FUNTYPE void CCNN_GEMM_FP16(fp16 *A_d, fp16 *B_d, fp16 *C_dOff, int offset, int m, int k, int n, int kGroupNum, int ifCAddin);
#endif
#endif

#ifdef _CMEM_USE_CPU_
void AvePoolComputerYKernel_FP16(fp16 *Y, fp16 *X, int N, int C, int H, int W, int kernelSize, \
                            int pType, int rtPad, int rbPad, int clPad, int crPad, int mainGlobalThreadId);
#else
#if (defined(_USE_CUDA_) || defined(_USE_ROCM_))
FUNTYPE void AvePoolComputerYKernel_FP16(fp16 *Y, fp16 *X, int N, int C, int H, int W, int kernelSize, \
                                    int pType, int rtPad, int rbPad, int clPad, int crPad);
#endif
#endif

#ifdef _CMEM_USE_CPU_
void AvePoolNCHWComputeDKernel_FP16(fp16 *pD, fp16 *D, int N, int C, int H, int W, \
                               int kernelSize, int pType, int rtPad, int rbPad, int clPad, int crPad, int mainGlobalThreadId);
#else
#if (defined(_USE_CUDA_) || defined(_USE_ROCM_))
FUNTYPE void AvePoolNCHWComputeDKernel_FP16(fp16 *pD, fp16 *D, int N, int C, int H, int W, \
                                       int kernelSize, int pType, int rtPad, int rbPad, int clPad, int crPad);
#endif
#endif

#ifdef _CMEM_USE_CPU_
void MaxPoolYKernel_FP16(fp16 *Y, fp16 *X, fp16 *rMat, fp16 *cMat, \
                    int N, int C, int H, int W, int kernelSize, int pType, int rtPad, \
                    int rbPad, int clPad, int crPad, int mainGlobalThreadId);
#else
#if (defined(_USE_CUDA_) || defined(_USE_ROCM_))
FUNTYPE void MaxPoolYKernel_FP16(fp16 *Y, fp16 *X, fp16 *rMat, fp16 *cMat, \
                            int N, int C, int H, int W, int kernelSize, int pType, int rtPad, \
                            int rbPad, int clPad, int crPad);
#endif
#endif

#ifdef _CMEM_USE_CPU_
void MaxPoolDKernel_FP16(fp16 *pD, fp16 *D, fp16 *rMat, fp16 *cMat, \
                    int N, int C, int H, int W, int kernelSize, int pType, int rtPad, \
                    int rbPad, int clPad, int crPad, int mainGlobalThreadId);
#else
#if (defined(_USE_CUDA_) || defined(_USE_ROCM_))
FUNTYPE void MaxPoolDKernel_FP16(fp16 *pD, fp16 *D, fp16 *rMat, fp16 *cMat, \
                            int N, int C, int H, int W, int kernelSize, int pType, int rtPad, \
                            int rbPad, int clPad, int crPad);
#endif
#endif

#ifdef _CMEM_USE_GPU_
#if (defined(_USE_CUDA_) || defined(_USE_ROCM_))
FUNTYPE void LayerNormCompYKernel_FP16(fp16 *xN, fp16 *Y, fp16 *mE, fp16 *var, \
                                  fp16 *inY, fp16 *gamma, fp16 *beta, int N, int C, int H, int W);
#endif
#endif

#ifdef _CMEM_USE_GPU_
#if (defined(_USE_CUDA_) || defined(_USE_ROCM_))
FUNTYPE void ComputerLNDeltaOfVarKernel_FP16(fp16 *cD, fp16 *inX, fp16 *dps, \
                                        fp16 *dEx, fp16 *dG, int CNHW);
#endif
#endif

#ifdef _CMEM_USE_GPU_
#if (defined(_USE_CUDA_) || defined(_USE_ROCM_))
FUNTYPE void LayerNormCompDKernel_FP16(fp16 *pD, fp16 *inX, fp16 *cD, fp16 *dE, \
                                  fp16 *dV, fp16 *dov, fp16 *donx, fp16 *dG, int N, \
                                  int C, int H, int W);
#endif
#endif

#ifdef _CMEM_USE_GPU_
#if (defined(_USE_CUDA_) || defined(_USE_ROCM_))
FUNTYPE void ComputerDeltaOfVarKernel_FP16(fp16 *cD, fp16 *inX, fp16 *dps, \
                                      fp16 *dG, fp16 *dEx, int CNHW);
#endif
#endif

#ifdef _CMEM_USE_GPU_
#if (defined(_USE_CUDA_) || defined(_USE_ROCM_))
FUNTYPE void BNCompPereDKernel_FP16(fp16 *rD, fp16 *inX, fp16 *cD, fp16 *dG, \
                               fp16 *dE, fp16 *dV, fp16 *donx, fp16 *dov, int N, \
                               int C, int H, int W);
#endif
#endif

#ifdef _CMEM_USE_CPU_
void BNCompYInferenceGPUKernel_FP16(fp16 *rY, fp16 *inY, fp16 *dG, fp16 *dB, \
                               fp16 *dSaveE, fp16 *dSaveV, int C, int H, int W, int mainGlobalThreadId);
#else
#if (defined(_USE_CUDA_) || defined(_USE_ROCM_))
FUNTYPE void BNCompYInferenceGPUKernel_FP16(fp16 *rY, fp16 *inY, fp16 *dG, fp16 *dB, \
                                       fp16 *dSaveE, fp16 *dSaveV, int C, int H, int W);
#endif
#endif

#ifdef _CMEM_USE_CPU_
void BNCompYTrainGPUKernel_FP16(fp16 *rY, fp16 *xN, fp16 *inY, fp16 *dG, fp16 *dB, \
                           fp16 *dE, fp16 *dV, int N, int C, int H, int W, int mainGlobalThreadId);
#else
#if (defined(_USE_CUDA_) || defined(_USE_ROCM_))
FUNTYPE void BNCompYTrainGPUKernel_FP16(fp16 *rY, fp16 *xN, fp16 *inY, fp16 *dG, \
                                   fp16 *dB, fp16 *dE, fp16 *dV, int N, int C, int H, int W);
#endif
#endif

#ifdef _CMEM_USE_CPU_
void L2NormCompYKernel_FP16(fp16 *outMat, fp16 *inMat, fp16 *cNorm, fp16 *scale, int N, \
                       int C, int H, int W, int mainGlobalThreadId);
#else
#if (defined(_USE_CUDA_) || defined(_USE_ROCM_))
FUNTYPE void L2NormCompYKernel_FP16(fp16 *outMat, fp16 *inMat, fp16 *cNorm, fp16 *scale, \
                               int N, int C, int H, int W);
#endif
#endif

#ifdef _CMEM_USE_CPU_
void L2NormCompDKernel_FP16(fp16 *rD, fp16 *cD, fp16 *yNorm, fp16 *xxt, fp16 *ddoty, \
                       fp16 *scale, int N, int C, int H, int W, int mainGlobalThreadId);
#else
#if (defined(_USE_CUDA_) || defined(_USE_ROCM_))
FUNTYPE void L2NormCompDKernel_FP16(fp16 *rD, fp16 *cD, fp16 *yNorm, fp16 *xxt, \
                               fp16 *ddoty, fp16 *scale, int N, int C, int H, int W);
#endif
#endif

#ifdef _CMEM_USE_CPU_
void MatNHWCToNCHWKernel_FP16(fp16 *outMatOff, fp16 *inMatOff, int off1, int off2, \
                         int N, int H, int W, int C, int mainGlobalThreadId);
#else
#if (defined(_USE_CUDA_) || defined(_USE_ROCM_))
FUNTYPE void MatNHWCToNCHWKernel_FP16(fp16 *outMatOff, fp16 *inMatOff, int off1, \
                                 int off2, int N, int H, int W, int C);
#endif
#endif

#ifdef _CMEM_USE_GPU_
#if (defined(_USE_CUDA_) || defined(_USE_ROCM_))
FUNTYPE void Mat4DAddinKernel_FP16(fp16 *dMat, fp16 *sMat, int N, int C, int H, int W);
#endif
#endif

#ifdef _CMEM_USE_CPU_
void MatNCHWToNHWCKernel_FP16(fp16 *outMatOff, fp16 *inMatOff, int off1, int off2, \
                         int N, int C, int H, int W, int mainGlobalThreadId);
#else
#if (defined(_USE_CUDA_) || defined(_USE_ROCM_))
FUNTYPE void MatNCHWToNHWCKernel_FP16(fp16 *outMatOff, fp16 *inMatOff, int off1, \
                                 int off2, int N, int C, int H, int W);
#endif
#endif

#ifdef _CMEM_USE_CPU_
void MatNCHWToCNHWKernel_FP16(fp16 *outMat, fp16 *inMat, int N, int C, int H, int W, int mainGlobalThreadId);
#else
#if (defined(_USE_CUDA_) || defined(_USE_ROCM_))
FUNTYPE void MatNCHWToCNHWKernel_FP16(fp16 *outMat, fp16 *inMat, int N, int C, int H, int W);
#endif
#endif

#ifdef _CMEM_USE_CPU_
void MatCNHWToNCHWKernel_FP16(fp16 *outMat, fp16 *inMat, int C, int N, int H, int W, int mainGlobalThreadId);
#else
#if (defined(_USE_CUDA_) || defined(_USE_ROCM_))
FUNTYPE void MatCNHWToNCHWKernel_FP16(fp16 *outMat, fp16 *inMat, int C, int N, int H, int W);
#endif
#endif

#ifdef _CMEM_USE_CPU_
void SumHWKernel_FP16(fp16 *y, fp16 *inY, int N, int C, int H, int W, int mainGlobalThreadId);
#else
#if (defined(_USE_CUDA_) || defined(_USE_ROCM_))
FUNTYPE void SumHWKernel_FP16(fp16 *y, fp16 *inY, int N, int C, int H, int W);
#endif
#endif

#ifdef _CMEM_USE_CPU_
void ChannleScaleCompYDKernel_FP16(fp16 *y, fp16 *inY, fp16 *scale, int N, int C, int H, \
                              int W, int mainGlobalThreadId);
#else
#if (defined(_USE_CUDA_) || defined(_USE_ROCM_))
FUNTYPE void ChannleScaleCompYDKernel_FP16(fp16 *y, fp16 *inY, fp16 *scale, \
                                      int N, int C, int H, int W);
#endif
#endif

#ifdef _CMEM_USE_CPU_
void BiasCompYKernel_FP16(fp16 *y, fp16 *inY, fp16 *bias, int N, int C, int H, int W, int mainGlobalThreadId);
#else
#if (defined(_USE_CUDA_) || defined(_USE_ROCM_))
FUNTYPE void BiasCompYKernel_FP16(fp16 *y, fp16 *inY, fp16 *bias, int N, int C, int H, int W);
#endif
#endif

#ifdef _CMEM_USE_CPU_
void ChangeSizeCompYKernel_FP16(fp16 *Y, fp16 *X, int N, int C, int H, int W, int cType, \
                           int tCS, int bCS, int lCS, int rCS, int mainGlobalThreadId);
#else
#if (defined(_USE_CUDA_) || defined(_USE_ROCM_))
FUNTYPE void ChangeSizeCompYKernel_FP16(fp16 *Y, fp16 *X, int N, int C, int H, int W, \
                                   int cType, int tCS, int bCS, int lCS, int rCS);
#endif
#endif

#ifdef _CMEM_USE_CPU_
void ChangeSizeCompDKernel_FP16(fp16 *pD, fp16 *D, int N, int C, int H, int W, int cType, \
                           int tCS, int bCS, int lCS, int rCS, int mainGlobalThreadId);
#else
#if (defined(_USE_CUDA_) || defined(_USE_ROCM_))
FUNTYPE void ChangeSizeCompDKernel_FP16(fp16 *pD, fp16 *D, int N, int C, int H, int W, \
                                   int cType, int tCS, int bCS, int lCS, int rCS);
#endif
#endif

#ifdef _CMEM_USE_GPU_
#if (defined(_USE_CUDA_) || defined(_USE_ROCM_))
FUNTYPE void SoftMaxImgCompYKernel_FP16(fp16 *maxMat, fp16 *maxPabMat, fp16 *Y, fp16 *inY, int N, int C, \
                                   int H, int W, int HW, int CHW, int subHWMatNum, int subWMatNum);
#endif
#endif

#ifdef _CMEM_USE_CPU_
void SoftMaxImgCompDKernel_FP16(fp16 *mLost, fp16 *pD, fp16 *iLab, fp16 *cY, \
                           int N, int C, int H, int W, int lostType, int mainGlobalThreadId);
#else
#if (defined(_USE_CUDA_) || defined(_USE_ROCM_))
FUNTYPE void SoftMaxImgCompDKernel_FP16(fp16 *mLost, fp16 *pD, fp16 *iLab, fp16 *cY, \
                                   int N, int C, int H, int W, int lostType);
#endif
#endif

#ifdef _CMEM_USE_CPU_
void SoftMaxCompDKernel_FP16(fp16 *resD, fp16 *inD, fp16 *inY, int bNum, int len, int mainGlobalThreadId);
#else
#if (defined(_USE_CUDA_) || defined(_USE_ROCM_))
FUNTYPE void SoftMaxCompDKernel_FP16(fp16 *resD, fp16 *inD, fp16 *inY, int bNum, int len);
#endif
#endif

#ifdef _CMEM_USE_CPU_
void DropoutLayerCompYKernel_FP16(fp16 *Y, fp16 *inY, fp16 *neuronPresent, fp16 *keepProb, int N, \
                             int C, int H, int W, int method, int mainGlobalThreadId);
#else
#if (defined(_USE_CUDA_) || defined(_USE_ROCM_))
FUNTYPE void DropoutLayerCompYKernel_FP16(fp16 *Y, fp16 *inY, fp16 *neuronPresent, fp16 *keepProb, \
                                     int N, int C, int H, int W, int method);
#endif
#endif

#ifdef _CMEM_USE_CPU_
void DropoutLayerCompDKernel_FP16(fp16 *pD, fp16 *cD, fp16 *neuronPresent, int N, int C, \
                             int H, int W, int mainGlobalThreadId);
#else
#if (defined(_USE_CUDA_) || defined(_USE_ROCM_))
FUNTYPE void DropoutLayerCompDKernel_FP16(fp16 *pD, fp16 *cD, fp16 *neuronPresent, int N, \
                                     int C, int H, int W);
#endif
#endif

#ifdef _CMEM_USE_CPU_
void PositionEncodeCompYKernel_FP16(fp16 *Y, fp16 *inY, int seqNum, int bNum, int wordDim, int mainGlobalThreadId);
#else
#if (defined(_USE_CUDA_) || defined(_USE_ROCM_))
FUNTYPE void PositionEncodeCompYKernel_FP16(fp16 *Y, fp16 *inY, int seqNum, int bNum, int wordDim);
#endif
#endif

#ifdef _CMEM_USE_CPU_
void FullConnectCompYKernel_FP16(fp16 *Y, fp16 *inY, fp16 *wData, int bNum, \
                            int inputNum, int outputNum, int mainGlobalThreadId);
#else
#if (defined(_USE_CUDA_) || defined(_USE_ROCM_))
FUNTYPE void FullConnectCompYKernel_FP16(fp16 *Y, fp16 *inY, fp16 *wData, \
                                    int bNum, int inputNum, int outputNum);
#endif
#endif

#ifdef _CMEM_USE_CPU_
void FullConnectCompDKernel_FP16(fp16 *dIn, fp16 *cD, fp16 *wData, int bNum, \
                            int inputNum, int outputNum, int mainGlobalThreadId);
#else
#if (defined(_USE_CUDA_) || defined(_USE_ROCM_))
FUNTYPE void FullConnectCompDKernel_FP16(fp16 *dIn, fp16 *cD, fp16 *wData, \
                                    int bNum, int inputNum, int outputNum);
#endif
#endif

#ifdef _CMEM_USE_CPU_
void SGDMOpt_Kernel_FP16(fp16 *w, fp16 *dm, fp16 *dw, fp16 *l2, fp16 *alpha, \
                    fp16 *ln, int len, int thisIndex);
#else
#if (defined(_USE_CUDA_) || defined(_USE_ROCM_))
FUNTYPE void SGDMOpt_Kernel_FP16(fp16 *w, fp16 *dm, fp16 *dw, fp16 *l2, \
                            fp16 *alpha, fp16 *ln, int len);
#endif
#endif

#ifdef _CMEM_USE_CPU_
void FullConnectCompWKernel_FP16(fp16 *tmpSumW, fp16 *yIn, fp16 *cD, int bNum, int inputNum, int outputNum, \
                            int mainGlobalThreadId);
#else
#if (defined(_USE_CUDA_) || defined(_USE_ROCM_))
FUNTYPE void FullConnectCompWKernel_FP16(fp16 *tmpSumW, fp16 *yIn, fp16 *cD, int bNum, int inputNum, int outputNum);
#endif
#endif

#ifdef _CMEM_USE_GPU_
#if (defined(_USE_CUDA_) || defined(_USE_ROCM_))
FUNTYPE void ReductionVectorSumKernel_FP16(fp16 *dA, fp16 *d_partial_sum, int len);
#endif
#endif

#ifdef _CMEM_USE_GPU_
#if (defined(_USE_CUDA_) || defined(_USE_ROCM_))
FUNTYPE void ReductionChannelsDotKernel_FP16(fp16 *dA, fp16 *dB, fp16 *d_partial_sum, int CNHW);
#endif
#endif

#ifdef _CMEM_USE_GPU_
#if (defined(_USE_CUDA_) || defined(_USE_ROCM_))
FUNTYPE void ReductionChannelsVarianceKernel_FP16(fp16 *dA, fp16 *d_partial_sum, fp16 *dMathExpect, int CNHW);
#endif
#endif

#ifdef _CMEM_USE_GPU_
#if (defined(_USE_CUDA_) || defined(_USE_ROCM_))
FUNTYPE void ReductionChannelsSumKernel_FP16(fp16 *dA, fp16 *d_partial_sum, int CNHW);
#endif
#endif

__host__ __device__ fp16 GetConvPlat2DData_FP16(int dstRIndex, int dstCIndex, int heightThisMat, int widthThisMat, int resultThisMatRow, \
                                           int resultThisMatCol, fp16 *thisChnImg, int height, int width, int kSize, \
                                           int pad, int stride, int dia, int ifFill, int fillStride);

#ifdef _CMEM_USE_CPU_
void GetMatByKernel4DRowInColOutKernel_FP16(fp16 *resultMat, fp16 *deviceK, int inC, int outC, int kSize, \
                                       int ifR180, int mainGlobalThreadId);
#else
#if (defined(_USE_CUDA_) || defined(_USE_ROCM_))
FUNTYPE void GetMatByKernel4DRowInColOutKernel_FP16(fp16 *resultMat, fp16 *deviceK, int inC, int outC, \
                                               int kSize, int ifR180);
#endif
#endif

#ifdef _CMEM_USE_CPU_
void Im2ColCompYDKKernel_FP16(fp16 *resultMat, fp16 *matOff, int offset, int H, int W, int height, int width, \
                         int kSize, int pad, int stride, int dia, int outStandPad, int heightThisMat, \
                         int widthThisMat, int resultThisMatRow, int resultThisMatCol, int matSize, int ifFill, \
                         int fillStride, int type, int mainGlobalThreadId);
#else
#if (defined(_USE_CUDA_) || defined(_USE_ROCM_))
FUNTYPE void Im2ColCompYDKKernel_FP16(fp16 *resultMat, fp16 *matOff, int offset, int H, int W, int height, int width, \
                                 int kSize, int pad, int stride, int dia, int outStandPad, int heightThisMat, \
                                 int widthThisMat, int resultThisMatRow, int resultThisMatCol, int matSize, \
                                 int ifFill, int fillStride, int type);
#endif
#endif

#ifdef _CMEM_USE_CPU_
void FormatToInCOutCKMat_Kernel_FP16(fp16 *gSum, fp16 *kUCMem, int inBatNum, int kSize, \
                                int inC, int outC, int mainGlobalThreadId);
#else
#if (defined(_USE_CUDA_) || defined(_USE_ROCM_))
FUNTYPE void FormatToInCOutCKMat_Kernel_FP16(fp16 *gSum, fp16 *kUCMem, int inBatNum, \
                                        int kSize, int inC, int outC);
#endif
#endif


#ifdef _CMEM_USE_GPU_
#if (defined(_USE_CUDA_) || defined(_USE_ROCM_))
FUNTYPE void SSDFirstRule_Kernel_FP16(fp16 *defaultBoxBuff, fp16 *tmpIou, fp16 *tmpIdx, fp16 *bx, fp16 *by, fp16 *bw, fp16 *bh, int flagV, int defaultBoxNum, \
                                 int boxStep, int flagI, int dbxI, int dbyI, int dbwI, int dbhI);
#endif
#endif

#ifdef _CMEM_USE_GPU_
#if (defined(_USE_CUDA_) || defined(_USE_ROCM_))
FUNTYPE void SSDSecondRule_Kernel_FP16(fp16 *defaultBoxBuff, fp16 *gBoxBuff, int flagV, int flagI, int matchIouI, int matchGbI, int defaultBoxNum, int dBoxStep, int dbxI, \
                                  int dbyI, int dbwI, int dbhI, int gNum, int gBoxStep, int gxI, int gyI, int gwI, int ghI);
#endif
#endif

#ifdef _CMEM_USE_GPU_
#if (defined(_USE_CUDA_) || defined(_USE_ROCM_))
FUNTYPE void ComputeNegativeSample_Kernel_FP16(fp16 *defaultBoxBuff, fp16 *Y, fp16 *resultArr, \
                                          fp16 *indexArr, fp16 *coff, int mLostType, int flagV, \
                                          int defaultBoxNum, int dBoxStep, int flagI, int offsetToPredictBox, int classNum);
#endif
#endif

#endif
