CUDA Event Timing과 Benchmark

cudabenchmarkprofilingoptimization

최적화의 첫 단계는 “빠른 것 같다”가 아니라 “같은 구간을 반복해서 재면 더 빠르다”이다.

host setup H2D copy kernel D2H copy verify
CPU timer
cudaEvent elapsed inside stream
PyTorch benchmark loop
최적화는 먼저 같은 구간을 같은 방식으로 재는 것에서 시작한다.

무엇을 잴 것인가

CPU timer로 전체 시간을 재면 host setup, memory copy, synchronization이 섞인다. kernel 자체 시간을 보려면 보통 CUDA event를 같은 stream에 기록한다.

cudaEventRecord(start);
kernel<<<grid, block>>>(...);
cudaEventRecord(stop);
cudaEventSynchronize(stop);
cudaEventElapsedTime(&ms, start, stop);

Colab에서의 현실

Colab에서는 Nsight Compute hardware counter 접근이 제한될 수 있다. 그래도 CUDA event, PyTorch torch.utils.benchmark, warmup loop만으로도 학습용 비교는 충분히 가능하다.

확인

  • CPU timer와 CUDA event가 재는 구간은 어떻게 다른가?
  • warmup을 하지 않으면 어떤 값이 섞일 수 있는가?
  • 최적화 전후 비교에서 input shape를 고정해야 하는 이유는 무엇인가?