#ifndef MATFP8_E4M3_KERNEL_BODY_H_
#define MATFP8_E4M3_KERNEL_BODY_H_

__host__ __device__ fp8_e4m3 GetSoftMaxGrad_FP8_E4M3(int lostType, int cIndex, int trueChannelIndex, fp8_e4m3 P_t, fp8_e4m3 P_i, float coff);
__host__ __device__ fp8_e4m3 GetSoftMaxLoss_FP8_E4M3(int lostType, fp8_e4m3 P_t, float coff);
__host__ __device__ int VectorMaxIndex_FP8_E4M3(fp8_e4m3 *arrayDataSrc, int veclength);
__host__ __device__ void LogSumExp_FP8_E4M3(fp8_e4m3 *result, fp8_e4m3 *arrFloat, int len, int tIndex);
__host__ __device__ void GetSoftMaxLossBySrcOut_FP8_E4M3(fp8_e4m3 *result, int lostType, fp8_e4m3 *arrFloat, int len, int tIndex, fp8_e4m3 coff);

#ifdef _CMEM_USE_GPU_
__device__ void AtomicFloat4AddIn_FP8_E4M3(fp8_e4m3 *dst, fp8_e4m3 *src);
#endif

__host__ __device__ union FP8_E4M3_vector4 FP8_E4M3_FLOAT4_Multi(union FP8_E4M3_vector4 f1, fp8_e4m3 coff);
__host__ __device__ union FP8_E4M3_vector4 FP8_E4M3_FLOAT4_ADD_Multi(union FP8_E4M3_vector4 f, union FP8_E4M3_vector4 f1, fp8_e4m3 coff);
   
__host__ __device__ fp8_e4m3 ActiveFun_FP8_E4M3(fp8_e4m3 x, int aType);

__host__ __device__ fp8_e4m3 DeActiveFun_FP8_E4M3(fp8_e4m3 y, int aType);

__host__ __device__ void MaxUpSampleYKernelFun_FP8_E4M3(_GMEM_ fp8_e4m3 *Y, _GMEM_ fp8_e4m3 *X, _GMEM_ fp8_e4m3 *rMat, _GMEM_ fp8_e4m3 *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_FP8_E4M3(_GMEM_ fp8_e4m3 *Y, _GMEM_ fp8_e4m3 *X, _GMEM_ fp8_e4m3 *rMat, _GMEM_ fp8_e4m3 *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_FP8_E4M3(_GMEM_ fp8_e4m3 *rD, _GMEM_ fp8_e4m3 *D, _GMEM_ fp8_e4m3 *rMat, _GMEM_ fp8_e4m3 *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_FP8_E4M3(_GMEM_ fp8_e4m3 *rD, _GMEM_ fp8_e4m3 *D, _GMEM_ fp8_e4m3 *rMat, _GMEM_ fp8_e4m3 *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_FP8_E4M3(_GMEM_ fp8_e4m3 *Y, _GMEM_ fp8_e4m3 *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, fp8_e4m3 f);

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

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

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

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

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

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

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

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

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

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

#ifdef _CMEM_USE_GPU_
__device__ void FloatBuffSet4ByteFloat_FP8_E4M3(fp8_e4m3 *dst, fp8_e4m3 *src);
#endif

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

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

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

#ifdef _CMEM_USE_GPU_
#if (defined(_USE_CUDA_) || defined(_USE_ROCM_))
__device__ void SetFloat4Data_FP8_E4M3(union FP8_E4M3_vector4 *dst, fp8_e4m3 *src);
#endif
#endif

#ifdef _CMEM_USE_GPU_
__device__ union FP8_E4M3_vector4 GetFloat4Data_FP8_E4M3(fp8_e4m3 *src);
#endif

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

#endif
