[MPC/CUDA] Error Check

간단한 CUDA kernel 예제와 함께 에러 처리에 대해 알아보자.

---

🏷 add_kernel 예제

SIZE가 8이고 1차원 A 배열에 1.0을 더한 값을 B 배열에 넣는 커널 함수이다.

//CUDA Kernel function
__global__ void add_kernel(float* b, const float* a){
    int i = threadIdx.x;
    b[i] = a[i] + 1.0f;
}

나중에 host의 main함수에서 add_kernel<<<1,SIZE>>>(dev_b, dev_a)를 통해 kernel call을 할 것이다. 

 

main함수는 다음과 같다. 

int main(void){
    const int SIZE = 8;
    const float a[SIZE] = {0.,1.,2.,3.,4.,5.,6.,7.,}
    float b[SIZE] = {0.};
    
    printf("..."); // 배열 a의 값들 출력
    
    float* dev_a = nullptr;
    float* dev_b = nullptr;
    
    cudaMalloc((void**)&dev_a, SIZE*sizeof(float));
    cudaMalloc((void**)&dev_b, SIZE*sizeof(float));
    
    cudaMemcpy(dev_a, a, SIZE*sizeof(float), cudaMemcpyHostToDevice); //dev_a = a;
    
    add_kernel<<<1,SIZE>>>(dev_b, dev_a);
    cudaDeviceSynchronize();
    
    cudaMemcpy(b, dev_b, SIZE*sizeof(float), cudaMemcpyDeviceToHost); // b = dev_b;
    
    printf("..."); // 배열 b의 값들 출력 
    
    cudaFree(dev_a);
    cudaFree(dev_b);
    return 0;
}

📂 CUDA function Rules

모든 CUDA 함수는 "cuda"로 시작한다. 예를 들어, cudaMalloc, cudaMemcpy... 등이 있다.

주의할 점은 CUDA function과 CUDA kernel function은 다른 개념이다!

또 대부분은 에러 코드를 리턴한다. 성공 시에는 cudaSuccess를 리턴한다. CUDA kernel은 에러 코드를 리턴하지 못하기 때문에 저번 포스팅에서 봤던 것처럼 cudaPeekAtLastError() 함수를 이용했다. 

간단한 예제를 보자.

if( cudaMalloc(&devptr, SIZE) != cudaSuccess){
    exit(1);
}

📄 cudaError_t

cudaError_t는 데이터 타입의 하나로, enum으로 typedef 되어있다. 즉, CUDA 함수에서 에러 코드를 반환할 때, String형식이 아닌 Number(숫자) 형식으로 리턴해주는 것을 의미한다.

가능한 value들은 다음과 같다.

이들은 모두 "cudaError + 이유" 와 같은 형식이다. 

 

그러면 에러 체크를 편하게 하기 위해 도와주는 함수들을 몇 가지 알아보자.

📄 cudaGetErrorName(err)

const char* cudaGetErrorName(cudaError_t err);

에러 코드는 enum value이기 때문에 바로 프린트를 실행하면 숫자 값이 나오게 된다. 따라서 해당하는 에러 코드의 이름을 반환해주는 함수이다.

Error code가 아닐 경우, NULL/nullptr을 반환한다. → 실패할 경우

📄 cudaGetErrorString(err)

const char* cudaGetErrorString(cudaError_t err);

이 함수는 에러의 이름이 아니라 그 에러에 대한 설명을 리턴해준다. 즉, 에러 발생의 원인을 말해주는 것이다. 

📂 cudaGetLastError()

이 함수는 이전 포스팅에서 알아봤으니, 중요한 특징 한 가지만 알아보자.

cudaGetLastError() 함수는 같은 host thread 안에서 발생한 에러에 대해서만 반환을 해준다. 다음 그림을 보자.

위 그림은 현재 3개의 Thread가 돌고 있는 상태이다. 이런 경우에 CUDA error 처리는 CPU thread 기준으로 한다는 것이다. thread1이 thread2에서 발생한 에러까지 관여하지 않는 것이다. 즉, CPU Thread 마다 별도의 error 상태를 관리한다.

즉, cudaGetLastError도 같은 CPU Thread 내에서 발생한 에러만을 알려주고, 다른 곳(thread)에서 발생한 에러에 대해서는 신경 쓰지 않는다. 

 

그렇다면, 맨 위에서 봤던 간단한 예제에 에러 체크 구문을 추가해보자.

...
    printf("..."); // 배열 b의 값들 출력
    
    cudaError_t err = cudaGetLastError();
    if(cudaSuccess != err){
        printf("CUDA:ERROR:cuda failure \"%s\"\n", cudaGetErrorString(err));
        exit(1);
    }
    else{
        printf("CUDA Success\n");
    }
    
    return 0;
}

여기서는 이전 포스팅과 다르게 에러를 처리하고 Error flag 값을 리셋시켜주기 위해 cudaPeekAtLastError 대신 cudaGetLastError함수를 사용했다. 그리고 에러가 발생할 경우, 에러의 원인을 설명하고 프로그램을 종료하는 방법으로 에러 처리를 수행했다.

📂 CUDA error check MACRO

이 에러 처리 코드를 매번 일일이 쓰지 말고, 매크로로 만들어 간편하게 사용해보자.

MPC의 모든 포스팅에 간혹, #include "./common.cpp" 이 등장하는데, 이와 같이 만든 매크로를 모두 common.cpp파일에 넣어두고 사용할 것이다.

🏷 CUDA_CHECK_ERROR()

#define CUDA_CHECK_ERROR() do{\
    cudaError_t e = cudaGetLastError();\
    if(cudaSuccess != e){\
        printf("CUDA failure \"%s\" at %s:%d\n",\
        cudaGetErrorString(e), __FILE__, __LINE__);\
        exit(1);\
    }\
}while(0)

여기서 do while(0)를 쓴 이유는, 보다시피 while(0) 뒤에 세미콜론(;)이 없다. 이 때문에, 코드 안에서 CUDA_CHECK_ERROR(); 와 같이 매크로에 세미콜론을 붙여야 돌아가게끔 만들기 위함이다.

__FILE__은 파일명, __LINE__은 라인 넘버를 반환해준다.

 

이를 이용해서 디버그 모드(Debug mode)와 릴리즈 모드(Release mode)에서 동작하는 매크로를 달리 할 수 있다.

대부분의 C/C++ 컴파일러에는 debug mode를 위한 _DEBUG와 release mode를 위한 NDEBUG가 정의되어 있다.

#if define(NDEBUG) // release mode
#define CUDA_CHECK_ERROR() 0
#else // debug mode
#define CUDA_CHECK_ERROR() do{\
    cudaError_t e = cudaGetLastError();\
    if(cudaSuccess != e){\
        printf("CUDA failure \"%s\" at %s:%d\n",\
        cudaGetErrorString(e), __FILE__, __LINE__);\
        exit(1);\
    }\
}while(0)
#endif

NDEBUG로 정의되어 있으면 release mode이므로 0;을 통해 아무것도 수행하지 않는다(무시). 반면, _DEBUG로 정의되어 있다면, CUDA Error를 체크하고 화면에 친절하게 에러의 원인을 print 해준다. __FILE__과 __LINE__은 C/C++ 컴파일러가 전체 프로그램을 컴파일하는 중에 자동으로 바꿔준다. 

하지만, 우리는 실제로는 if문을 제거하고 항상 error를 체크하는 방법을 사용할 것이다. 이러한 방법이 있다는 것만 알아두자.

 

그럼 다음과 같이 다시 코드를 변경할 수 있다.

printf("...") // 배열 b의 값들 출력

//error check
CUDA_CHECK_ERROR();

return 0;
}

 

제대로 동작하는지 확인하기 위해, 고의로 에러를 발생시켜 보자. 위 코드에서 Device에서 계산한 결과 값을 Host의 B배열로 복사할 때, 마지막 파라미터를 오류가 나도록 변경한다.

...
cudaMemcpy(b, dev_b, SIZE*sizeof(float), cudaMemcpyDeviceToDevice); // ERROR!!
...
CUDA_CHECK_ERROR();

return 0;
}

정상적인 코드라면 cudaMemcpyDeviceToHost를 수행해야 한다.

이 코드를 실행해보면 다음과 같은 결과를 얻을 수 있다.

b는 host(main)의 주소 값이다. 하지만, DeviceToDevice를 수행했기 때문에, 참조할 수 없는 주소 값이 되어 "invailid argument"라는 에러가 발생한 것을 확인할 수 있다. 당연히 b 배열에 결과 값 복사에도 실패했을 것이다. 뒤에 __FILE__, __LINE__에 따라 파일명과 라인 넘버가 출력된 것도 확인할 수 있다.