CUDA Event를 이용한 커널(SAXPY) 실행시간 측정

호스트의 타이머 함수를 사용할 경우, cudaDeviceSynchronize( )와 같은 명시적인 배리어 함수를 사용해야한다. 이 경우 GPU 파이프라인이 stall되는 문제 등이 발생될 수 있다. 따라서 CUDA 런타임에서 제공하는 event API를 사용해서 커널의 실행 시간을 측정하는 방식을 이용하는 것이 좋다.

CUDA 이벤트 데이터 타입인 cudaEvent_t 를 사용하는 다음의 함수들을 사용한다.

• __host__ cudaError_t cudaEventCreate( cudaEvent_t* event)
• __host__ __device__ cudaError_t cudaEventDestroy(cudaEvent_t* event)
• __host__ __device__ cudaError_t cudaEventRecord (cudaEvent_t* event cudaStream_t stream=0)
• __host__ cudaError_t cudaEventElapsedTime( float* ms, cudaEvent_t start, cudaEvent_t end)

전체적인 사용방식은 다음과 같다.

cudaEventCreate( )를 통해서 시작과 끝을 기록할 start 이벤트와 stop 이벤트를 생성한다.

커널이 시작되기 직전에 cudaEventRecord( )를 사용하여 start 이벤트에 시작시점을 기록한다.

마찬가지로 커널이 끝나는 지점에 cudaEventRecord( )를 사용하여 stop 이벤트에 종료시점을 기록한다. 

이 때 호스트 코드에는 cudaEventSynchronize( )를 이용하여 stop 이벤트가 기록될 때까지, 다음 CPU 실행을 블록한다.

마지막으로 cudaEventElapsedTime( )을 이용해서 두 이벤트 사이의 시간을 구한다.

( cudaEventDestroy( )를 통해서 두 이벤트를 삭제해준다. )

 

참고로 CUDA Runtime API 문서에서는 이벤트 생성은 cudaEventCreate( ) 이외에도 cudaEventCreateWithFlags( )로도 가능하다고 되어 있다. 그리고 cudaEventSynchronize( )는 cudaEventBlockingSync 플래그를 가진 이벤트에 대해서는 CPU 스레드를 디바이스가 이벤트를 종료시킬 때까지 블록한다고 하고, 만약 이 플래그가 설정되지 않은 경우는 CPU 스레드는 busy waiting 한다고 되어 있다. 

일단 cudaEventCreate( )가 디폴트로 어떤 플래그로 생성되는지 이 문서에 기록되어 있지는 않지만,  Mark Harris의 문서에 따르면 단순히 블록한다고 나와있기 때문에 cudaEventBlockingSync( )로 설정되지 않을까하고 추정해본다. 그리고 단 하나의 커널에 대한 시간 측정을 주로 해서인지 cudaEventSynchronize( )를 사용하지 않는 예제들도 많았다.

 

https://hayunjong83.tistory.com/20

에서 작성하였던 SAXPY 코드에 CUDA Event를 이용한 시간측정을 추가해보았다. 

간단한 연산이었고, 5회 반복측정 평균 0.01301 ms 정도의 실행시간을 보였다.

 

참고)

https://devblogs.nvidia.com/how-implement-performance-metrics-cuda-cc/

https://docs.nvidia.com/cuda/cuda-runtime-api/group__CUDART__EVENT.html#group__CUDART__EVENT