V100에 추가된 Tensor Core라는 것은 무엇인가 ?

Tesla V100 GPU에 새로 추가된 Tensor Core라는 것은 기존의 Tesla P100 GPU에는 없던 것입니다.  이 Tensor Core라는 것은 무엇이고 어떤 일에 사용할 수 있으며, 그로 인한 성능 향상은 어느 정도일까요 ?

1. Tensor Core라는 것은 무엇인가 ?

NVIDIA의 사전적 설명에 따르면, Tensor Core는 matrix-multiply-and-accumulate (행렬 곱셈+덧셈) 장치로서 Tensor Core 1개당 하나의 사이클에서 64개의 floating point FMA mixed-precision operation을 수행할 수 있습니다.  FMA란 fused multiply–add로서, 곱셈과 덧셈을 하나의 operation에서 한꺼번에 수행하는 것입니다.  V100 GPU에는 SM당 8개씩의 Tensor Core가 장착되어 있으므로, 하나의 SM에서는 64 x 8 x 2(곱셈 + 덧셈) = 1024번의 floating point 연산을 하나의 사이클에서 수행가능합니다.  V100에는 80개의 SM이 장착되어 있으므로 한 사이클에 80 * 1024 = 8196번의 floating point 연산이 가능합니다.

Tensor Core의 특징은 저 위에서 말한 대로, deep learning에서 많이 사용되는 D = A * B + C와 같은 행렬의 곱셈과 덧셈을 한번에 수행해버린다는 점입니다.  그 외에도 매우 중요한 점이 있습니다.  바로 mixed precision 연산이 매우 쉽게, 그것도 한 사이클에서 가능하다는 것입니다.  즉, D = A * B + C의 연산에서, A와 B는 FP16(floating point 16-bit, half-precision)으로, C와 D는 FP32(floating point 32-bit, single-precision)으로 한번에 처리하는 것이 가능합니다.

2. Tensor Core는 어떤 일에 사용되나 ?

한줄 요약하면 deep learning에서 빠르면서도 높은 정확도(accuracy)의 모델을 만들 때 매우 유용하게 사용됩니다.

Deep learning training phase에서 하는 일은 수십~수백만장의 이미지에서 공통되는 feature를 추출하는 것이므로, input data에 굳이 FP32를 사용할 필요가 없고 FP16 정도면 충분한 경우가 많습니다.  그럴 경우, P100 Pascal GPU부터는 FP32 장치 하나로 FP16 연산 2개를 수행할 수 있으므로 성능도 2배로 높일 수 있고, 특히 GPU 메모리를 절약할 수 있습니다.  그러나 deep learning의 목표는 빠른 training이 아니라 더 높은 정확도(accuracy)입니다.  더 높은 accuracy를 위해서는, 그렇게 FP16의 행렬 data를 곱한 결과를 합산할 때 FP32로 해야 합니다.  즉, deep learning에서는 이런 mixed precision의 곱셈과 덧셈이 매우 자주 사용됩니다.  이런 경우에 Tensor Core를 사용하는 것이 특히 좋습니다.


3.  Tensor Core는 어떻게 사용하면 되는가 ?

물론 기존의 application을 그대로 V100에서 수행한다고 해서 Tensor Core를 자동으로 쓰는 것은 아닙니다.   SW가 Tensor Core를 사용하도록 개발되어야 하지요.  2018년 3월인 지금 현재로서는 CUDA library 중 cuBLAS와 cuDNN이 Tensor Core를 지원합니다.  cuBLAS와 cuDNN은 모두 caffe와 tensorflow 등 주요 deep learning framework에서 모두 사용하는 library입니다.  cuBLAS는 GEMM(행렬끼리의 곱셈) 연산을 수행할 때, 그리고 cuDNN은 CNN 및 RNN을 가속하는데 Tensor Core를 사용합니다.

그렇다고 cuBLAS 및 cuDNN을 쓰는 모든 application이 Tensor Core에 의한 성능 가속 효과를 보는 것은 아닙니다.  가령 다음과 같이 cuBLAS를 사용할 때 Tensor Core를 쓰도록 application을 수정해야 합니다.

// Set the math mode to allow cuBLAS to use Tensor Cores:
cublasStat = cublasSetMathMode(handle, CUBLAS_TENSOR_OP_MATH);
...
// Invoke the GEMM, ensuring k, lda, ldb, and ldc are all multiples of 8, 
// and m is a multiple of 4:
cublasStat = cublasGemmEx(handle, transa, transb, m, n, k, alpha,
                          A, CUDA_R_16F, lda,
                          B, CUDA_R_16F, ldb,
                          beta, C, CUDA_R_16F, ldc, CUDA_R_32F, algo);

cuBLAS에서 GEMM 연산을 한다고 해도, 위에서처럼 math mode가 CUBLAS_TENSOR_OP_MATH로 맞춰져 있어야 하고, k, lda, ldb, ldc 등이 모두 8의 배수여야 합니다.  이런 조건에 맞지 않으면 GEMM 연산은 Tensor Core를 이용하지 않고 기존 CUDA core를 이용합니다.

현재는 tensorflow, pytorch, caffe2 등의 framework이 Tensor Core의 혜택을 보도록 적용이 되어 있습니다.


4. Tensor Core로 인한 성능 향상은 어느 정도인가 ?

GPU에는 single-precision(FP32) 연산장치와 double-precision(FP16) 연산장치가 각각 따로 있으며, V100에는 그 비율이 2:1로 되어 있습니다.   그래서 V100의 FP32 성능은 15.7 TFLOPS인데, FP64 성능은 7.8 TFLOPS라고 하는 것입니다.  여태까지는 보통 deep learning에서는 (더 높은 정확도를 위해서) FP32를 사용하고, 과학기술 연산이나 금융 연산에서는 FP64를 사용합니다.  그런데 NVIDIA의 자료에 V100의 "Deep Learning" 성능은 125 TFLOPS라고 되어 있습니다.  이건 실제 deep learning 성능이라기보다는, tensor operation 성능을 말하는 것입니다.  위에서 V100은 한 사이클에 8196번의 floating point 연산이 가능하다고 했고 V100의 GPU boost clock speed가 1.53GHz이므로 8196 * 1530 million = 125 TFLOPS가 나오는 것입니다.

NVIDIA에서 내놓은 자료에 따르면 단순 GEMM 연산의 경우 V100에서 Tensor Core를 이용하면 기존 P100 대비 4~9배의 성능을 냅니다.

역시 NVIDIA에서 내놓은 자료에 따르면 cuDNN convolution 속도가 V100에서 Tensor Core를 이용하면 기존 P100 대비 4~5배 빨라집니다.

그러나 실제로 tensorflow를 이용하여 training을 해보면 언제나 그렇게 빨라지는 것은 아닙니다.  일단 deep learning 전체 과정이 모두 GEMM으로 이루어진 것도 아니요 cuDNN convolution으로만 된 것도 아니기 때문입니다.   가령 아래의 tensorflow github에 올라온 issue를 보십시요.

Tensor Core support for NVIDIA Volta architecture #15897

이 URL의 내용을 요약하면 "Volta GPU에서 FP16으로 tensorflow를 수행할 때, NVIDIA의 광고에 따르면 tensor core 덕분에 4배~8배의 성능이 나와야 하지만 실제로는 2배의 성능만 나오더라, 어떻게 된 것인가 ?"라는 이슈 제기에 대해, tensorflow 개발측이 "tensorflow는 어떤 내부 알고리즘을 사용하는 것이 가장 좋은 성능을 낼지 찾아내는 자동 튜닝 장치를 갖추고 있고, 여기에 이미 tensor core를 사용하도록 적용이 되어 있다, 그러나 자동 튜닝 결과가 언제나 항상 tensor core를 사용하도록 나오는 것은 아니다"라고 답한 것입니다.  특히 이 질문 속 트레이닝의 경우 NCHW에서 NHWC로의 이미지 포맷 전환에서 많은 시간을 소비하므로 특히 그런 것으로 나옵니다. 

* 용어 해설
NCHW :  Num_samples x Channels x Height x Width
NHWC : Num_samples x Height x Width x Channels


윗 글의 대부분은 아래 URL에 나온 내용을 알기 쉽게 풀어 놓은 것이며 성능 자료들의 출처도 아래 URL입니다.

https://devblogs.nvidia.com/programming-tensor-cores-cuda-9/