nVIDIA cudnnConvolutionBackwardFilter를 조사하는 루프 라인 모델

TL;DR



cudnnConvolutionBackwardFilter() 을 1000가지 파라미터로 프로파일링하고
GTX1060의 지붕 라인 모델에 플롯했다.



많은 경우에, 100Gflops 이상의 성능이 나오는 한편, 극단적으로 성능 열화하는 파라미터가 존재하는 것을 알 수 있었다.
CUDNN_CONVOLUTION_BWD_FILTER_ALGO_FFT도 일단 사용되는 것이 있음을 알 수 있었다.

평가 대상




기계
nVIDIA GeForce GTX1060 3G


피크 연산 성능
3.469Tflops

피크 밴드 폭
192GB/sec

CUDA
v10.0

cuDNN
v7.4


평가 조건 (파라미터)



1000가지의 파라미터를 랜덤하게 생성. 가격 범위는 아래 표와 같습니다.


매개변수
가치 영역


배치 크기 n
randint(1, 32)

입력 채널 ci
randint(1, 20)

입력 이미지 크기 Hi
randint(1, 122)

입력 이미지 크기 Wi
randint(1, 122)

스트라이드 u (h 방향)
randint(1, 7)

커널 kernel_h
randint(1, 7)

커널 kernel_w
randint(1, 7)

패딩 pad_h
randint(0, 8)

패딩 pad_w
randint(0, 8)

dilation_h
1 고정

dilation_w
1 고정


※ 알고리즘은 1000가지의 파라미터 각각에 대해 가장 빠른 알고리즘을 선택하고 있다.
cudnnGetConvolutionBackwardFilterAlgorithm(..., CUDNN_CONVOLUTION_BWD_FILTER_PREFER_FASTEST,...)

평가식


  • Performance [Gflops] = (연산 횟수)/(실행 시간)
  • Arithmetic intensity [GB/sec] = (연산 횟수)/(입출력 데이터 사이즈)
  • (연산 횟수): c++에서 convolution backward filter를 어리석게 써서 세었다 . mul과 add는 2operation으로 계산됩니다.
  • (실행 시간) : cudaEventElapsedTime()
  • (입출력 데이터 사이즈) = (n*ci*hi*wi + n*co*ho*wo + co*ci*kernel_h*kernel_w)*4 [Byte]


  • 평가 환경



    cudnnConvolutionBackwardFilter()


    ※ Cython의 setpu.py만은 여러가지 하드 코딩하고 있으므로 리포지토리에 넣지 않습니다. 죄송합니다.

    참고문헌


  • 지붕 라인 모델 정보 :
  • 소스 코드 세트

  • Cython으로 c++ Warp 방법에 대해 :
  • ㅡㅡㅡㅜㅜㅜㅜ jp / st 레아 민 g / gpg 푸 / 아 d ゔ 센세 d_gpg 뿌 / 2015 / 아 d ゔ 퓌 세 d_gpg 푸 06. pdf

  • htps : // 기주 b. m / PyS rm / pys rm / uki / St rin gs-an dby by s-in-Cy

  • CUDA Runtime에 의한 실행 시간 측정 정보 :
  • htps : // 메이 m. 코 m / @ 유스 켄 / 캇 ぃ ー ー ー ー ー ー ー ー ー ー ー ー ー ーcb461b6d8

  • 좋은 웹페이지 즐겨찾기