#ifndef MATFP32_KERNEL_BODY_H_
#define MATFP32_KERNEL_BODY_H_

__host__ __device__ fp32 GetSoftMaxGrad_FP32(int lostType, int cIndex, int trueChannelIndex, fp32 P_t, fp32 P_i, float coff);
__host__ __device__ fp32 GetSoftMaxLoss_FP32(int lostType, fp32 P_t, float coff);
__host__ __device__ int VectorMaxIndex_FP32(fp32 *arrayDataSrc, int veclength);
__host__ __device__ void LogSumExp_FP32(fp32 *result, fp32 *arrFloat, int len, int tIndex);
__host__ __device__ void GetSoftMaxLossBySrcOut_FP32(fp32 *result, int lostType, fp32 *arrFloat, int len, int tIndex, fp32 coff);

#ifdef _CMEM_USE_GPU_
__device__ void AtomicFloat4AddIn_FP32(fp32 *dst, fp32 *src);
#endif

__host__ __device__ union FP32_vector4 FP32_FLOAT4_Multi(union FP32_vector4 f1, fp32 coff);
__host__ __device__ union FP32_vector4 FP32_FLOAT4_ADD_Multi(union FP32_vector4 f, union FP32_vector4 f1, fp32 coff);
   
__host__ __device__ fp32 ActiveFun_FP32(fp32 x, int aType);

__host__ __device__ fp32 DeActiveFun_FP32(fp32 y, int aType);

__host__ __device__ void MaxUpSampleYKernelFun_FP32(_GMEM_ fp32 *Y, _GMEM_ fp32 *X, _GMEM_ fp32 *rMat, _GMEM_ fp32 *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_FP32(_GMEM_ fp32 *Y, _GMEM_ fp32 *X, _GMEM_ fp32 *rMat, _GMEM_ fp32 *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_FP32(_GMEM_ fp32 *rD, _GMEM_ fp32 *D, _GMEM_ fp32 *rMat, _GMEM_ fp32 *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_FP32(_GMEM_ fp32 *rD, _GMEM_ fp32 *D, _GMEM_ fp32 *rMat, _GMEM_ fp32 *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_FP32(_GMEM_ fp32 *Y, _GMEM_ fp32 *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, fp32 f);

__host__ __device__ void MatrixAvePooling_FP32(_GMEM_ fp32 *Y, _GMEM_ fp32 *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, fp32 f);

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

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

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

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

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

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

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

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

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

#ifdef _CMEM_USE_GPU_
__device__ void FloatBuffSet4ByteFloat_FP32(fp32 *dst, fp32 *src);
#endif

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

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

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

#ifdef _CMEM_USE_GPU_
#if (defined(_USE_CUDA_) || defined(_USE_ROCM_))
__device__ void SetFloat4Data_FP32(union FP32_vector4 *dst, fp32 *src);
#endif
#endif

#ifdef _CMEM_USE_GPU_
__device__ union FP32_vector4 GetFloat4Data_FP32(fp32 *src);
#endif

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

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

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

#ifdef _CMEM_USE_CPU_
void AvePoolComputerYKernel_FP32(fp32 *Y, fp32 *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_FP32(fp32 *Y, fp32 *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_FP32(fp32 *pD, fp32 *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_FP32(fp32 *pD, fp32 *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_FP32(fp32 *Y, fp32 *X, fp32 *rMat, fp32 *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_FP32(fp32 *Y, fp32 *X, fp32 *rMat, fp32 *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_FP32(fp32 *pD, fp32 *D, fp32 *rMat, fp32 *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_FP32(fp32 *pD, fp32 *D, fp32 *rMat, fp32 *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_FP32(fp32 *xN, fp32 *Y, fp32 *mE, fp32 *var, \
                                  fp32 *inY, fp32 *gamma, fp32 *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_FP32(fp32 *cD, fp32 *inX, fp32 *dps, \
                                        fp32 *dEx, fp32 *dG, int CNHW);
#endif
#endif

#ifdef _CMEM_USE_GPU_
#if (defined(_USE_CUDA_) || defined(_USE_ROCM_))
FUNTYPE void LayerNormCompDKernel_FP32(fp32 *pD, fp32 *inX, fp32 *cD, fp32 *dE, \
                                  fp32 *dV, fp32 *dov, fp32 *donx, fp32 *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_FP32(fp32 *cD, fp32 *inX, fp32 *dps, \
                                      fp32 *dG, fp32 *dEx, int CNHW);
#endif
#endif

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

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

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

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

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

#ifdef _CMEM_USE_CPU_
void MatNHWCToNCHWKernel_FP32(fp32 *outMatOff, fp32 *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_FP32(fp32 *outMatOff, fp32 *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_FP32(fp32 *dMat, fp32 *sMat, int N, int C, int H, int W);
#endif
#endif

#ifdef _CMEM_USE_CPU_
void MatNCHWToNHWCKernel_FP32(fp32 *outMatOff, fp32 *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_FP32(fp32 *outMatOff, fp32 *inMatOff, int off1, \
                                 int off2, int N, int C, int H, int W);
#endif
#endif

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

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

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

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

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

#ifdef _CMEM_USE_CPU_
void ChangeSizeCompYKernel_FP32(fp32 *Y, fp32 *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_FP32(fp32 *Y, fp32 *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_FP32(fp32 *pD, fp32 *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_FP32(fp32 *pD, fp32 *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_FP32(fp32 *maxMat, fp32 *maxPabMat, fp32 *Y, fp32 *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_FP32(fp32 *mLost, fp32 *pD, fp32 *iLab, fp32 *cY, \
                           int N, int C, int H, int W, int lostType, int mainGlobalThreadId);
#else
#if (defined(_USE_CUDA_) || defined(_USE_ROCM_))
FUNTYPE void SoftMaxImgCompDKernel_FP32(fp32 *mLost, fp32 *pD, fp32 *iLab, fp32 *cY, \
                                   int N, int C, int H, int W, int lostType);
#endif
#endif

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

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

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

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

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

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

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

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

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

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

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

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

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

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

#ifdef _CMEM_USE_CPU_
void Im2ColCompYDKKernel_FP32(fp32 *resultMat, fp32 *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_FP32(fp32 *resultMat, fp32 *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_FP32(fp32 *gSum, fp32 *kUCMem, int inBatNum, int kSize, \
                                int inC, int outC, int mainGlobalThreadId);
#else
#if (defined(_USE_CUDA_) || defined(_USE_ROCM_))
FUNTYPE void FormatToInCOutCKMat_Kernel_FP32(fp32 *gSum, fp32 *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_FP32(fp32 *defaultBoxBuff, fp32 *tmpIou, fp32 *tmpIdx, fp32 *bx, fp32 *by, fp32 *bw, fp32 *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_FP32(fp32 *defaultBoxBuff, fp32 *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_FP32(fp32 *defaultBoxBuff, fp32 *Y, fp32 *resultArr, \
                                          fp32 *indexArr, fp32 *coff, int mLostType, int flagV, \
                                          int defaultBoxNum, int dBoxStep, int flagI, int offsetToPredictBox, int classNum);
#endif
#endif

#endif
