#ifndef PADDLE_FLUID_PLATFORM_ERR_MSG_CU
#define PADDLE_FLUID_PLATFORM_ERR_MSG_CU

#include"paddle/fluid/platform/enforce.h"
#define __LEN_ERROR_MSG 100

typedef  struct _CudaKerenlErrorPro{
    char _CudaErrorMsg[__LEN_ERROR_MSG];
    int _line;
    char _file[__LEN_ERROR_MSG];
    }CudaKerenlErrorPro;

    #define PADDLE_ENFORCE_CHECK_CUDA_KERNEL()                    \
    do {                                                          \
        char msg[__LEN_ERROR_MSG]={0};                            \
        int line=0;                                               \
        char occur;                                               \
        char file[__LEN_ERROR_MSG]={0};                           \
        get_msg_from_cuda(&occur,msg,file,&line);                 \
        if(occur) {                                               \
        throw ::paddle::platform::EnforceNotMet(msg,file,line);   \
      }                                                           \
    } while (0)

    #define PADDLE_ENFORCE_CUDA_KERNEL(CON, MSG)             \
    do {                                                     \
      if (!(CON)) {                                          \
    send_error_msg(MSG,__FILE__,__LINE__);                   \
                         asm("exit;");                       \
      }                                                      \
    } while (0)

  __device__ static char _CudaKernelErrorOccurred[1]={0};
  __device__ static CudaKerenlErrorPro _CudaKernelErrorMsg;
  
  __device__  __host__ inline void _strcpy(char* src,char* target){
      int i=0;
      while(0!=src[i] && i<__LEN_ERROR_MSG-1){
          target[i]=src[i];
          i++;
      }
      target[i]=0;
  }
  
   __device__ __host__ inline void send_error_msg(char*msg,char* file,int line){
      if(!_CudaKernelErrorOccurred[0]){
          _strcpy(msg,_CudaKernelErrorMsg._CudaErrorMsg);
          _strcpy(file,_CudaKernelErrorMsg._file);
          _CudaKernelErrorMsg._line=line;
          _CudaKernelErrorOccurred[0]=1;
      }
  }

extern "C"{
 void  inline get_msg_from_cuda(char* con,char* msg,char* file,int*line){
        CudaKerenlErrorPro temp;
        char occur[1]={0};
        cudaError_t err=cudaMemcpyFromSymbol(&temp,(_CudaKernelErrorMsg), sizeof(temp));
        cudaMemcpyFromSymbol(occur,_CudaKernelErrorOccurred,sizeof(char));
        strcpy(msg,temp._CudaErrorMsg);
        strcpy(file,temp._file);
        *line=temp._line;
        *con=occur[0];
        PADDLE_ENFORCE_CUDA_SUCCESS(cudaGetLastError());
    }
}

#endif