CUDA

CUDA는 NVIDIA 사에서 제공하는 GPGPU 솔루션(CUDA API) 및 그 실행 주체가 되는 그래픽 카드 내의 연산부(CUDA Core)를 뜻한다. 현재 GPGPU를 잘 활용하는 신경망 인공지능 연구 분야에서 절대적인 입지를 가지고 있다

탄생 배경과 현황

그래픽 카드는 원래 2D 및 3D 렌더링을 위해 대량의 SIMD 명령어를 위주로 한꺼번에 작동하도록 특별히 설계된 IC 칩이기 때문에 CPU의 벡터 명령어 집합처럼 벡터 프로세싱에 특화되어 있다. 이때 컴퓨터공학 전문가들이 GPU를 가지고 과학적 연구를 위한 수치 연산을 시도할 것이라는 것을 예상한 NVIDIA는 2006년 발표된 G80 GPU(Geforce 8000번대)부터 MIMD 명령어도 집어넣어 최초로 그래픽스 연산이 아닌 벡터 프로세싱 연산도 지원하도록 개선하고 관련 툴 프로그램(SDK)도 제공하였다. 이에 뒤늦게 2009년에 Apple이 주도하여 개발한 OpenCL을 중심으로 인텔과 AMD가 반격을 시도했으나 이미 3년이나 벌어진 발표 시기 사이에 CUDA를 활용한 각종 라이브러리가 쏟아져 나왔고 OpenCL의 기능은 CUDA도 다 지원할 수 있었기 때문에 이 CUDA 라이브러리들을 전부 OpenCL로 포팅하는 시도는 대부분 좌절되었고 현재까지도 CUDA가 OpenCL에 비해 범용 연산에 대한 수요 측면에서 우위를 차지하는 환경이 조성되었다. 이후로 나오는 ROCm이나 OneAPI 같은 후발주자도 가뿐하게 물리치는 상황이다.

TensorFlow, Pytorch 같은 신경망 AI 프레임워크들도 모두 CUDA를 사용한 GPU 기반 AI 가속을 중점적으로 지원한다

사용법

우선 사용자가 C/C++ 프로그래밍 언어를 알고 있다는 전제 하에 CPU를 Host, GPU를 Device로 서술한다.

1. 먼저 드라이버와 함께 nvcc라는 프로그램이 컴퓨터에 설치되어 있는지 확인한다. nvcc는 CUDA를 사용하는 프로그램 내부의 CUDA 명령어를 활용하는 함수(Kernel이라고 한다)를 먼저 컴파일해주는 전처리를 담당한다.
2. CUDA 프로그램 소스 코드의 확장자는 일반적으로 cu이다. 따라서 example.cu 같은 파일을 생성한다.
3. 1.에서 다룬 Kernel을 작성한다. 이때 함수 정의 앞에 __global__ 같이 작성할 함수가 Kernel이라는 표시를 꼭 붙인다.
4. 이제 Kernel을 사용하는 CPU 측 코드를 짜되 다음과 같은 순서로 작성한다.

1. 먼저 dim3 같은 데이터 형으로 GPGPU 기능을 사용하기 위한 Grid와 Block을 정의한다. 여기서 Grid는 Block 꾸러미이고, Block은 실제 Kernel의 병렬 처리 실행 단위인 Thread의 꾸러미라고 생각하자. 가령 Geforce에서는 하나의 Grid에 Block이 최대 32개까지 존재하며 하나의 Block 안에 1024개의 Thread가 존재하므로 이 중에 필요한 만큼 Grid와 Block 사이즈를 정한다.
2. cudamalloc 함수를 통해 그래픽 카드의 VRAM의 일부를 kernel에 넘길 데이터 할당 공간으로 받아온다.
3. cudamemcpy 함수에 파라미터를 VRAM 상의 데이터 공간, Host 프로그램의 데이터 공간, 복사할 데이터 크기, 데이터 복사 방향(cudaHostToDevice, cudaDeviceToHost 등) 순서로 넘긴다.
4. Kernel을 호출한다.
5. Kernel 호출이 끝나면 cudafree로 VRAM 상의 데이터 공간을 할당 해제해야 한다.

• Kernel을 호출할 때에는 KernelFunc<<<Grid, Block>>>(param1, ...); 식으로 호출한다.
• Makefile 등으로 프로그램 컴파일 순서를 지정할 때 nvcc 컴파일러->일반 컴파일러 순서로 작동하도록 작성하여 프로그램을 기계어로 바꾼다.
• 주의할 점으로, Kernel 내부에서는 if~else 구문을 사용하지 않는 것이 강력히 권장된다. 만일 Kernel 내부에 if문 같은 조건문이 나올 경우 nvcc는 GPU의 연산 자원을 반으로 갈라 조건문이 참일 경우와 거짓일 경우를 모두 계산한 뒤 맞는 연산만 취사 선택하는 식으로 비효율적인 계산을 하게 된다. 조건문이나 반복문은 CPU 상에서 실행되는 프로그램에서 처리해야 한다.

TMI

• 일반적으로 게임에서 활용하는 32bit 부동 소수점 연산(단정밀도 연산)은 NVIDIA의 모든 카드가 같은 GPU 클럭에서 동일한 성능을 제공하나 그래픽스 작업(영상 편집, 렌더링)에서 주로 활용하는 64bit 부동소수점 연산(배정밀도 연산)은 게이머들이 많이 사가는 Geforce보다 그래픽스 전문가용 그래픽 카드인 Quadro 및 데이터 센터에서의 수치 연산 전용 그래픽 카드인 Tesla 시리즈에서 더 빠르다. 이는 Geforce 400 시리즈에서 NVIDIA가 절대 성능 우위를 사수하기 위해 무리하게 CUDA 코어를 때려 박았다가 엄청난 발열로 혹평을 듣자 급히 하위 제품군에서 배정밀도 연산부를 크게 쳐냈기 때문이다. 그런데 이게 전화위복이 되어 Geforce 700번대부터 Geforce 10 시리즈까지 NVIDIA 그래픽 카드의 전력 소모 대비 성능 향상 폭이 매우 크게 되어 배정밀도 연산을 포기할 수 없었던 AMD의 GCN 아키텍처 기반 라데온을 Geforce가 저발열 고성능 특성으로 찍어 누르는 원동력이 되었다.