이번은 호스트 프로그램으로 Kernel나 CUDA API의 실행 시간을
측정하는 방법에 대해 설명합니다. Kernel 내부에서의 시간측정에 관해서는 CUDA 프로그램의 시간을
측정하는 - clock()을 참고하세요.
시간측정의
주의점!
CUDA API는 그 실행 문맥 중에서 제일 처음에 불린 API는 매우
시간이 걸리는 특징이 있습니다.이것은 CUDA의
초기화 작업을 하고 있기 때문입니다. 그러므로, cudaSetDevice등의 다른 초기화 함수를 부르는지와 최초의 API는 시간에 필요한 시간을 잘 살펴봐야
합니다.
시간측정을
통일합니다.
CUDA로 프로그램을 구현하여 프로그램의 성능을 측정할 필요가 있을 때
시간을 측정하는 방법으로 고민하는 일이
있습니다. OS에 의한 시간측정방법과 완전히 다르기 때문에 portability 코드를 쓰는 것은 상당히 귀찮습니다. 잘못된 시간측정을 실시하면 결과는 터무니없는 값이
되겠지만, 계측 시간이 msec 단위인 경우 그것이 작성된 쓴 원시 코드 때문인지 아니면 정말로 그
시간 걸렸는지 알 수 없습니다.
이번에는 CUDA의 Event를 사용하여 어떤 환경에도 이용할 수 있고
단순하고 올바른 측정 시간을 얻을 수
있습니다.
사용예제
#include
<stdio.h>
#define
VEC_NUM (512)
__global__
void dummy(float *dA)
{
int id =
blockDim.x * blockIdx.x + threadIdx.x;
dA[id] =
dA[id] + dA[VEC_NUM-1-id];
}
int main(int
argc, char** argv)
{
float
A[VEC_NUM];
float
*dA;
size_t
size = VEC_NUM * sizeof(float);
int
i;
for(
i=0; i<VEC_NUM; i++ ) {
A[i]
= (float) i;
}
cudaMalloc( (void **) &dA, size
);
cudaMemcpy( dA, A, size,
cudaMemcpyHostToDevice );
float
elapsed_time_ms=0.0f;
cudaEvent_t start,
stop;
cudaEventCreate( &start
);
cudaEventCreate( &stop
);
cudaEventRecord( start, 0
);
dummy<<<VEC_NUM/128,
128>>>(dA);
cudaEventRecord( stop, 0
);
cudaEventSynchronize( stop
);
cudaEventElapsedTime( &elapsed_time_ms,
start, stop );
cudaMemcpy( A, dA, size,
cudaMemcpyDeviceToHost );
printf(
"time: %8.2f ms\n", elapsed_time_ms );
cudaEventDestroy( start
);
cudaEventDestroy( stop
);
cudaFree( dA
);
return
0;
}
프로그램의 내용에 큰
의미는 없습니다. 중요한 것은 다음 부분입니다.
float
elapsed_time_ms=0.0f;
cudaEvent_t start,
stop;
cudaEventCreate( &start
);
cudaEventCreate( &stop
);
cudaEventRecord( start, 0
);
dummy<<<VEC_NUM/128,
128>>>(dA);
cudaEventRecord( stop, 0
);
cudaEventSynchronize( stop
);
cudaEventElapsedTime( &elapsed_time_ms,
start, stop );
CUDA Event를 사용하여 CUDA API에서 비동기화로 실행되고
있는 여러가지 명령(호스트·디바이스 사이의
데이터 전송이나, kernel의 실행등)의 진행 상황을 체크 포인트(이벤트) 마다 모니터 할 수 있습니다.
1.cudaEvent_t형의 변수를 기록하고 싶은 이벤트의 수 만큼 선언
한다
2.cudaEventCreate를 이용해 이벤트를
생성한다
3.cudaEventRecord를 이용해 이벤트를
기록한다
4.cudaEventSynchronize로 이벤트를 동기화
한다
5.cudaEventElapsedTime로 이벤트간의 시간측정을
한다
6.cudaEventDestroy로 이벤트를
파기한다
cudaEventRecord는 다음과 같은 동작을
합니다.
이 함수가 불리기
직전까지 호출된 CUDA API의 동작이 모두 끝난 직후에 이벤트를
기록한다.
덧붙여 이 함수는 비동기화에 동작하기 때문에 경과시간을 계측하기 전에
이벤트를 기록하고 있는지를 보증할 필요가 있습니다. 이를 위한 것이
cudaEventSynchronize입니다. 함수 인수의 이벤트가 종료 할 때까지 cudaEventSynchronize는 실행을
블록합니다.
마지막으로,
cudaEventElapsedTime로 시간을 측정합니다. 단위는 밀리
세컨드입니다.
![/!\](http://www.viper.pe.kr/wiki/sinorca4moin/img/alert.png "/!\"){kind=small}
코드는 준비중 - 
준비되면 업로드 예정
No comments:
Post a Comment