IBM이 자랑하는 POWER9 + V100 GPU의 조합이 H2O DriverlessAI에 어떤 도움이 될까 ?

H2O Driverless AI는 model training 중 GPU에서 돌리는 것이 더 효과적이라고 판단될 때마다 필요에 따라 드문드문 GPU를 사용합니다.   이는 우측 하단의 "GPU usage" tab을 클릭하면 아래와 같이 보실 수 있습니다.

물론 nvidia-smi 에서도 모니터링하실 수 있습니다.   GPU 1개당 1개의 process가 수행되는 것이 아니라, GPU 메모리가 허용하는 한 필요에 따라 GPU 1개당 여러 개의 process들이 수행되기도 합니다.

Fri Oct  5 03:36:45 2018
+-----------------------------------------------------------------------------+
| NVIDIA-SMI 396.26                 Driver Version: 396.26                    |
|-------------------------------+----------------------+----------------------+
| GPU  Name        Persistence-M| Bus-Id        Disp.A | Volatile Uncorr. ECC |
| Fan  Temp  Perf  Pwr:Usage/Cap|         Memory-Usage | GPU-Util  Compute M. |
|===============================+======================+======================|
|   0  Tesla P100-SXM2...  On   | 00000002:01:00.0 Off |                    0 |
| N/A   34C    P0    64W / 300W |    455MiB / 16280MiB |     40%      Default |
+-------------------------------+----------------------+----------------------+
|   1  Tesla P100-SXM2...  On   | 00000003:01:00.0 Off |                    0 |
| N/A   36C    P0    77W / 300W |    455MiB / 16280MiB |     38%      Default |
+-------------------------------+----------------------+----------------------+
|   2  Tesla P100-SXM2...  On   | 0000000A:01:00.0 Off |                    0 |
| N/A   32C    P0    71W / 300W |    455MiB / 16280MiB |     40%      Default |
+-------------------------------+----------------------+----------------------+
|   3  Tesla P100-SXM2...  On   | 0000000B:01:00.0 Off |                    0 |
| N/A   36C    P0    64W / 300W |    455MiB / 16280MiB |     38%      Default |
+-------------------------------+----------------------+----------------------+

+-----------------------------------------------------------------------------+
| Processes:                                                       GPU Memory |
|  GPU       PID   Type   Process name                             Usage      |
|=============================================================================|
|    0     94004      C   ...el-running(prot=False)-XGBoostModel-fit   445MiB |
|    1     94011      C   ...el-running(prot=False)-XGBoostModel-fit   445MiB |
|    2     94044      C   ...el-running(prot=False)-XGBoostModel-fit   445MiB |
|    3     94126      C   ...el-running(prot=False)-XGBoostModel-fit   445MiB |
+-----------------------------------------------------------------------------+


과연 여기서 NVLink가 어느 정도의 효과를 발휘할까요 ?  그것을 확인하는 가장 쉽고 생생한 방법은 nvprof 명령을 이용하여 DriverlessAI process들이 GPU를 어떻게 사용하는지 profile을 떠보는 것입니다.

먼저, DAI의 홈 디렉토리에 있는 start-dai.sh를 다음과 같이 수정합니다.  원래의 line은 #으로 막아놓았고, 아래 붉은색의 명령어들이 새로 삽입된 nvprof 명령어입니다.

[root@p57a22 dai-1.3.1-linux-ppc64le]# cat start-dai.sh
#!/bin/bash

if [ "x${RUN_DAI_CREATE_PIDFILE}" == "x1" ]; then
# python -m h2oai &
 /usr/local/cuda-9.2/bin/nvprof --log-file /tmp/nvprof1.%p.txt --export-profile /tmp/nvprof1.%p.nvvp --print-gpu-trace --profile-child-processes python -m h2oai &
    echo $! > ${RUN_DAI_PIDFILE_DIR}/dai.pid
else
# exec python -m h2oai
    exec /usr/local/cuda-9.2/bin/nvprof --log-file /tmp/nvprof1.%p.txt --export-profile /tmp/nvprof1.%p.nvvp --print-gpu-trace --profile-child-processes python -m h2oai
fi


위와 같이 해놓고 kill-dai.sh를 수행하여 기존 DAI를 stop시킨 뒤, run-dai.sh를 수행하여 DAI를 새로 구동한 뒤에 큼직한 dataset으로 training을 해봅니다.   이때, 다음과 같이 child process마다 txt와 nvvp로 텍스트와 바이너리 로그 파일이 생성되는 것을 보실 수 있습니다.

[root@p57a22 dai-1.3.1-linux-ppc64le]# ls -l /tmp/nvprof1*
...
-rw-r--r-- 1 root root 170096640 Oct  5 03:16 /tmp/nvprof1.83368.nvvp
-rw-r--r-- 1 root root 233180816 Oct  5 03:22 /tmp/nvprof1.83368.txt
-rw-r--r-- 1 root root 180073472 Oct  5 03:16 /tmp/nvprof1.83405.nvvp
-rw-r--r-- 1 root root 246428557 Oct  5 03:20 /tmp/nvprof1.83405.txt
...

이 중 하나를 CUDA에 딸려 있는 NVIDIA Visual Profiler (NVVP)로 분석해보면 다음과 같습니다.

"CUDA memcpy HtoD" (Host to Device, 즉 서버 CPU에서 GPU로의 memcpy도 9 GB/s에 가깝고 DtoD (GPU간의 memcpy)는 134 GB/s에 가깝게 사용되는 것을 보실 수 있습니다.

특히 8개의 GPU를 장착한 일반 x86 기반 GPU 서버들은 NVLink를 1개씩만 이용하여 연결되므로 GPU간의 통신은 이론상 50 GB/s, CPU와 GPU간의 통신은 이론상 32 GB/s가 최대인 것에 비해, IBM AC922은 3개의 NVLink를 aggregation시켜 CPU와 GPU이나 GPU끼리나 모두 150 GB/s로 연결된 것이 H2O DriverlessAI의 운용에 크게 유리하다는 것을 쉽게 보실 수 있습니다.

세계 최대의 수퍼컴 Summit에 대한 가벼운 이야기들

2018년 6월 9일, 드디어 세계 최대의 수퍼컴 Summit이 발표되었습니다.   요구되었던 '최소 150 PetaFLOPS 이상'을 훌쩍 뛰어넘은 200 PetaFLOPS을 구축했고, 이로서 중국의 TaihuLight에게 넘겨주었던 기존의 세계 최대 수퍼컴이라는 타이틀을 5년만에 미국이 되찾아 왔습니다.

수퍼컴이 사용되는 HPC (High Performance Computing)라는 분야는 과학기술 쪽의 박사님들 외에는 잘 모르시고 또 알 필요도 없는 분야입니다만, 이번 글에서는 IT가 아닌 분들도 쉽고 흥미있게 읽으실 수 있도록 이 세계 최대의 수퍼컴에 대해 써봤습니다.

(Summit에 관련된 인포그래픽입니다.   원본은 https://www.olcf.ornl.gov/wp-content/uploads/2018/06/Summit_bythenumbers_FIN.png 입니다.)


1.  CORAL project란 무엇인가 ?

전세계 수퍼컴의 연산능력은 1등부터 500등까지 순위대로 top500.org에 등재됩니다. 여기에 얼마나 많은 수퍼컴이 등재되었는지, 또 얼마나 큰 수퍼컴이 등재되는지가 월드컵 순위나 GDP 순위처럼 일종의 국력의 척도로서 자존심 싸움이 되기도 합니다.  그런데 2013년, 중국이 자체 개발한 TaihuLight라는 수퍼컴이 미국을 따돌리고 세계 최대의 수퍼컴이 되었습니다.  그 성능은 100 PetaFLOPS(Peta = 10의 15승, FLOPS = FLoating Point Operations Per Second = 초당 부동소수점 연산)였습니다.

꼭 이것 때문이라고 하기는 어렵겠으나, 미국의 수퍼컴 프로젝트를 주관하는 미국 에너지성(Department of Energy)에서는 바로 다음해인 2014년 CORAL (Collaboration of Oak Ridge, Argonne, and Livermore)이라는 이름의 프로젝트 하에 오크 릿지, 아르곤, 로렌스 리버모어의 세 국립 연구소의 수퍼컴을 업그레이드/대체하는 사업을 발표합니다.  그 중 오크 릿지와 로렌스 리버모어는 IBM와 NVIDIA, 그리고 인피니밴드 업체인 Mellanox의 OpenPOWER 컨소시엄이 수주한 것입니다.   아르곤에서는 인텔이 주사업자로서 Intel Xeon CPU와 Xeon Phi 가속기를 이용한 수퍼컴을 구축하게 되었습니다.   이 세 연구소의 수퍼컴 중 가장 큰 것은 Summit 입니다.


2.  오크 릿지 Summit 성능이 발표 되었으니, 이제 구축이 끝난 것인가 ?

아닙니다.  아직 Summit은 구축 중이며, 내년 초에나 정식 오픈하는 것으로 되어 있습니다.  물론 최근인 6월 8일, 오크 릿지는 Summit의 성능이 200 PetaFLOPS라고 발표했습니다.  원래 2014년에 나왔던 제안요청서(RFP)에는 150 PetaFLOPS 이상의 성능을 가질 것이 요구되었는데, 그 기준을 여유있게 통과한 것입니다.   다만 6월 20일 현재, 아직 top500.org에는 등재되지 않았습니다.  Top500 list는 매년 6월과 11월, 2차례에 걸쳐 등재되는데, 6월 말에나 업데이트가 될 예정입니다.


3.  최근 볼턴 미 백악관 안보 보좌관이 북한의 핵탄두를 오크 릿지로 반출해야 한다고 주장했는데, 그 오크 릿지가 이 오크 릿지인가 ?

맞습니다.  테네시 주의 오크 릿지에 위치한 오크 릿지 국립 연구소는 제2차 세계대전 당시 미국의 핵무기 개발을 위한 극비 프로젝트였던 맨하탄 계획이 수행된 곳입니다.  다만 볼턴 보좌관이 언급한 곳이 이 국립 연구소를 뜻하는 것인지 아니면 바로 옆에 있는 핵연료 처리 시설인 Y-12 National Security Complex를 가리키는 것인지는 불분명하다고 합니다.

(오크 릿지 연구소 바로 옆에 위치한 Y-12 핵시설입니다.  여기서 제2차 세계대전 때 히로시마와 나가사키에 투하된 원폭을 만들었습니다.)


4.  왜 같은 CORAL 프로젝트 중에서 오크 릿지 연구소의 Summit에 대해서만 이야기가 많고, 로렌스 리버모어 연구소의 Sierra에 대해서는 발표되는 것이 거의 없는가 ?

두가지 이유가 있습니다.  첫째, 오크 릿지의 Summit이 더 크고 더 먼저 시작했습니다.  둘째, 로렌스 리버모어의 보안이 훨씬 엄중합니다.

오크 릿지에서도 고에너지 물리학, 즉 핵물리학 연구를 계속 합니다만, 일반 민간 연구용으로도 컴퓨팅 파워를 빌려줍니다만, 로렌스 리버모어 연구소는 정말 순수하게 핵무기 등 국가 안보 관련된 연구를 하는 곳이라서 외부인들이 Sierra에 접속하여 민간용 연구를 할 일은 없을 것 같습니다.   가령 조지 클루니와 니콜 키드먼 주연의 1997년 영화 The Peacemaker에서, 니콜 키드먼이 백악관 소속 핵무기 전문가로 나오는데, 백악관에서 일하기 전에는 로렌스 리버모어에서 핵무기 개발을 했던 것으로 설정되어 있습니다.


5.  그렇다면 오크 릿지의 Summit은 일반인들도 사용이 가능한가 ?

가능은 합니다.  이 링크 https://www.olcf.ornl.gov/for-users/getting-started/#applying-for-a-user-account 를 보면 여러 가지 형태의 user account 신청이 가능합니다.  다만 이 곳은 아마존이나 MS Azure 같은 클라우드 장사하는 곳이 아니므로 '돈을 낼테니 쓰게 해달라'는 식의 요청은 안 되는 모양이고 연구 과제 등에 대해 까다로운 심사를 받아야 하는 모양입니다. 

최근 top500.org에 나온 기사 https://www.top500.org/news/openpower-gathers-momentum-with-major-deployments/ 를 보면 다음과 같이 우버도 자사의 tensorflow 기반의 신경망을 이 Summit에서 training하고자 오크 릿지와 협업 중에 있다고 합니다.

"우버는 텐서플로우에 기반한 자사의 분산 training 프레임웍인 Horovod를 Summit에서 수행하기를 원하고 있습니다.  이 딥러닝 업무에는 여러가지 우버의 앱들, 가령 자율주행 네비게이션, 여정 예측, 사기 방지 등이 포함됩니다.  우버는 특히 GPU 사용과 관련되어 Horovod의 확장성을 한단계 더 넓히는데 관심이 있습니다."


6.  Summit이라는 수퍼컴이라는 것은 어떻게 생긴 컴퓨터인가 ?

겉에서 보면 일반 데이터 센터에 있는 서버 랙들의 모양새와 별로 다를 것이 없습니다.

7.  Summit 같은 수퍼컴에도 모니터와 키보드가 달려 있는가 ?

극단적으로 한줄 요약하면 없습니다.  Summit은 4,608대의 AC922이라는 IBM GPU 서버를 모아서 만든 것이므로, 이런 수천대의 서버에 일일이 모니터와 키보드 마우스를 연결하여 제어하는 것은 무리입니다.  이런 서버에는 BMC라고 하는 이더넷 포트(RJ-45)처럼 생긴 BMC(Baseboard Management Controller)용 포트가 있는데, 이를 통해 IPMI (Intelligent Platform Management Interface)라는 것을 통해 서버를 제어합니다.

8. Summit에 사용되는 CPU는 어떤 것인가 ?

과거에는 vector processor라는, 수학 계산에 특화된 아키텍처의 프로세서가 수퍼컴의 주류를 이루었습니다만, 지금은 대부분 사라졌습니다.  수퍼컴에도 보통의 회사에서 사용하는 웹서버에 장착된 일반적인 프로세서가 주로 사용됩니다.  Summit에는 IBM의 POWER9이 사용되었습니다.  이 POWER9 프로세서도 엄청난 고가의 특수 프로세서는 아닙니다.  인터넷 쇼핑몰의 웹서버보다 훨씬 더 중요한 업무라고 할 수 있는 은행이나 증권사의 계정계 서버에는 RISC 프로세서인 IBM POWER 프로세서가 사용되는 경우가 많으며, 그런 곳에서 사용되는 것과 똑같은 프로세서입니다.


9. Summit에서는 어떤 OS를 사용하는가 ?

프로세서와 마찬가지입니다.  요즘은 수퍼컴에서도 윈도우즈도 많이 사용합니다.  물론 대세는 리눅스입니다.  이번에 Summit에 올라가는 OS도 POWER 계열 프로세서의 아키텍처인 ppc64le용 Redhat 7.5 Linux가 올라가 있습니다.


10. 똑같은 CPU와 똑같은 OS를 쓴다면, 수퍼컴은 대체 무엇이 특수한가 ?

초기의 수퍼컴은 위에서 언급한 것처럼 수학 계산용 특수 CPU를 매우 빠른 속도로 돌리는 방식이었습니다.  지금도 수퍼컴하면 흔히 생각나는 Cray가 바로 그런 수퍼컴의 선구자였지요.  그러나 무작정 CPU clock speed를 높은 GHz로 돌리는 것에는 반도체 물성에 따른 한계가 있고 또 일반 프로세서의 성능이 좋아지면서 점차 일반 프로세서의 서버들 수십~수천 대를 고성능 네트워크로 엮는 cluster 방식이 주종을 이루게 되었습니다.  이러한 cluster 방식의 수퍼컴은 그 클러스터를 이루는 개개의 서버 성능도 중요하지만, 그것을 연결하는 네트워크의 성능 및 병렬 프로그래밍을 위한 SW 기술도 중요합니다.  또한 전력을 적게 사용해야 하고, 많은 data를 고속으로 처리하기 위한 병렬 파일시스템이 필요합니다.

Summit은 IBM의 AC922이라는 2-socket 서버 4,608대로 이루어져 있는데, 이 서버들은 2장의 POWER9 프로세서 뿐만 아니라, 가속기로 NVIDIA의 최신 Telsa GPU인 V100을 6장씩 장착했습니다.  그리고 발열 문제를 효율적으로 해결하기 위해 수냉식으로 되어 있습니다.

11.  Summit은 과거의 수퍼컴과 무엇이 다른가 ?

가장 큰 차이는 GPU 기술을 본격적으로 채택한 수퍼컴이라는 점입니다.  Summit이 대체하게 되는 오크 릿지의 기존 수퍼컴인 Titian과 비교를 해보면 그 차이는 명확합니다.  Summit은 Titian보다 노드(node) 수, 즉 구성 요소인 서버 대수는 1/4 수준이면서도 성능은 7.4배를 냅니다.  덕분에 그렇게 성능을 많이 내면서도 전력소비량도 기존의 1.4배에 불과합니다.

(이 표에는 노드 당 성능이 42 TFLOPS라고 나와 있습니다만, 이는 double-precision, 즉 FP64 성능을 기준으로 한 것입니다.  또한 원래 V100 SXM2의 FP64 성능은 NVIDIA에 따르면 7.8 TFLOPS이므로, 7.8 x 6장 = 46.8 TFLOPS가 맞을텐데, Oak Ridge의 홈페이지에는 42 TFLOPS라고 나옵니다.  오타일까요 ? )


12. 결국 Summit의 핵심은 GPU인 것 같다.   요즘 Deep Learning에 GPU 서버를 많이 사용하는데, 그런 GPU 서버와 이 Summit은 무엇이 다른가 ?

기본적으로는 동일합니다.  그러나 중요한 차이가 2가지 있습니다.  하나는 NVLink이고, 다른 하나는 PCIe Gen4 입니다.  둘다 I/O 쪽 기술이지요.

NVLink는 NVIDIA가 개발한 GPU 연결 기술로서, 일반 x86 서버에서도 사용됩니다.  그러나 IBM POWER9에서만 가능한 것이 따로 있습니다.  PCIe Gen3 포트만 장착된 Intel Xeon 프로세서에서는 NVLink를 GPU간의 연결에만 사용할 수 있고, Intel Xeon 프로세서와 GPU 간의 통신은 느린 PCIe를 사용해야 합니다.  그와 달리, IBM POWER9에는 NVLink 2.0 port가 실리콘에 박혀 있습니다.  이를 통해 POWER9 프로세서는 PCIe가 아닌, NVLink를 통해 V100 GPU와 직접 통신을 합니다.  GPU 간의 peer-to-peer 통신이 NVLink로 되어 있을 경우 Caffe 등을 이용한 deep learning에서는 어느 정도 효과를 볼 수 있지만, HPC 업무에는 오로지 CPU와 GPU 간의 통신이 중요합니다.

또 있습니다.  IBM POWER9은 2018년 여름 현재, PCIe Gen4 포트가 장착된 유일한 프로세서입니다.  이를 통해 AC922은 고속 네트워크인 100Gb EDR Infiniband를 일반 x86 서버보다 2배의 대역폭으로 연결할 수 있습니다.  이 infiniband를 통해 Summit은 고속 병렬파일시스템을 연결하고, 또 Summit을 이루는 4,608대의 AC922 서버 간의 MPI (Message Passing Interface) 통신을 고속으로 수행합니다.

이러한 이유 때문에 IBM와 NVIDIA, 그리고 인피니밴드 업체인 Mellanox의 OpenPOWER 컨소시엄이  Summit과 Sierra라고 하는 2대의 수퍼컴을 개발 구매하는 CORAL project의 공급자로 채택된 것입니다.

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/