GPU 병렬 프로그래밍 CUDA의 모든 것

텐서플로우 외에도, 많은 다른 딥러닝 프레임워크가 Caff2, CNTK, Databricks, H2O.ai, Keras, Torcho, Torch를 포함한 GPU 지원을 위해 CUDA를 사용한다. 대부분의 경우 심층 신경망 컴퓨팅을 위해 cuDNN 라이브러리를 사용한다. 이 라이브러리는 딥러닝 프레임워크의 훈련에 매우 중요하기 때문에 특정 버전의 cuDNN을 사용하는 모든 프레임워크는 동등한 사용례에 대해 기본적으로 동일한 성능 수치를 갖는다. CUDA와 cuDNN이 버전마다 개선되면서 새로운 버전으로 업데이트하는 모든 딥러닝 프레임워크가 동일한 성능 상의 이득을 누리게 된다. 성능이 프레임워크마다 다른 이유는 여러 GPU 및 여러 노드로 얼마나 잘 확장하느냐에 있다.

CUDA 프로그래밍


CUDA 툴킷
CUDA 툴킷에는 라이브러리, 디버깅 및 최적화 도구, 컴파일러, 문서화 및 애플리케이션 배포를 위한 런타임 라이브러리가 포함된다. CUDA 툴킷은 딥러닝, 선형 대수, 신호 처리 및 병렬 알고리즘을 지원하는 요소들을 가지고 있다. 일반적으로 CUDA 라이브러리는 엔비디아 GPU의 모든 제품군을 지원하지만, 딥러닝 교육 워크로드의 경우 P100보다 3배 가량 빠른 V100과 같은 최신 모델에 가장 적합하다. 필요한 알고리즘이 적절한 라이브러리에 구현되어 있다면, 하나 이상의 라이브러리를 사용하는 것이 GPU를 가장 쉽게 활용할 수 있는 방법이다.



CUDA 딥러닝 라이브러리
딥러닝 영역에는 세 가지 주요 GPU 가속화 라이브러리가 존재한다. cuDNN은 앞서 말했듯 대부분의 오픈소스 딥러닝 프레임워크의 GPU 구성요소이며, TensorRT는 엔비디아의 고성능 딥러닝 추론 최적화기이자 런타임이라 할 수 있다. 또, 비디오 추론 라이브러리인 딥스트림(DeepStream)도 있다. TensorRT는 신경망 모델을 최적화하고, 낮은 정확도를 보정하며, 훈련된 모델을 클라우드, 데이터센터, 내장 시스템 또는 자동차 제품 플랫폼에 배포하기도 한다.

CUDA 선형 대수 및 수학 라이브러리
선형 대수는 텐서 연산과 딥러닝의 기반이 된다. 그 동안 과학자들과 엔지니어들은 1989년 포트란에서 구현된 매트릭스 알고리즘의 집합체인 BLAS(Basic Linear Algear Subprograms)를 사용해 왔다. cuBLAS는 BLAS의 GPU 가속화 버전이며, GPU를 사용해 가장 높은 성능의 연산을 수행하는 방법이다. cuBLAS는 매트릭스가 조밀하다고 가정한다. cuSPARSE가 매트릭스를 처리한다.

CUDA 신호 처리 라이브러리
FFT(Fast Fourier Transform)는 신호 처리에 사용되는 기본 알고리즘 중 하나이며, 신호(예를 들어 오디오 파형)를 주파수의 스펙트럼으로 변환하는 기능을 한다. cuFFT는 GPU 가속화된 FFT라 할 수 있다. H.264와 같은 표준을 사용하는 코덱은 전송 및 디스플레이를 위해 비디오를 인코딩/압축 및 디코딩/디컴프레스 한다. 엔비디아 비디오 코덱 SDK는 GPU를 통해 이 프로세스의 속도를 높여준다.

CUDA 병렬 알고리즘 라이브러리
병렬 알고리즘을 위한 세 라이브러리는 모두 다른 목적을 가지고 있다. NCCL (Nvidia Collective Communications Library)은 여러 GPU 및 노드에 걸쳐 앱을 확장하기 위한 것이며, nvGRAPH 는 병렬 그래프 분석을 위한 것이다. Thrust는 CUDA를 기반으로 하는 C++ 스탠다드 템플릿 라이브러리이다. Thrust는 스캔, 정렬 및 감소와 같은 다양한 데이터 병렬 원형을 제공한다.

CUDA vs. CPU 성능
경우에 따라 CPU 기능 대신 드롭인 형식의 CUDA 기능을 사용할 수도 있다. 예를 들어 BLAS의 GEMM 행렬 곱셈(matrix multiplication) 루틴은 NVBLAS 라이브러리에 연결하기만 하면 GPU 버전으로 대체할 수 있다.



CUDA 프로그래밍 기본
알맞은 CUDA 라이브러리 루틴을 찾을 수 없는 경우 저수준 CUDA 프로그래밍을 시도해 볼 수밖에 없다. 그래도 요즘은 2000년대 말보다는 훨씬 쉽게 할 수 있을 것이다. 여러 가지 이유가 있겠지만, 무엇보다 신택스(syntax)가 쉬워지고, 개발 툴도 더 나아졌기 때문이다. 필자의 유일한 불만이라면, 맥OS에서는 최신 CUDA 컴파일러와 최신 C++ 컴파일러(Xcode)가 동기화되는 경우가 거의 없다는 것이다. 때문에 애플에서 옛날 명령어줄 툴을 다운로드하고 xcode-select를 사용하여 해당 도구로 전환하여 CUDA 코드를 컴파일하고 연결해야 한다.

예를 들어, 다음 2개의 어레이를 추가하는 간단한 C/C++ 루틴을 보자.

void add(int n, float *x, float *y)
{
for (int i = 0; i < n; i++)
y[i] = x[i] + y[i];
}

선언에 __global__ 키워드를 추가해 GPU에서 실행될 커널로 전환하고, 삼중 괄호 구문을 사용하여 커널을 호출할 수 있다.

add<<<1, 1>>>(N, x, y);

또한 malloc/new와 free/delete 콜을 cudaMallocManaged와 cudaFree로 바꿔야만 GPU에서 공간을 할당할 수 있다. 마지막으로, CPU 결과를 사용하기 전에 GPU 계산이 완료될 때까지 기다려야 한다. 이것은 cudaDeviceSynchronize를 통해 구현할 수 있다.

삼중 괄호는 하나의 쓰레드 블록과 하나의 쓰레드를 사용한다. 현재 엔비디아 GPU는 많은 블록과 쓰레드를 처리할 수 있다. 예를 들어, 파스칼 GPU 아키텍처를 기반으로 하는 테슬라 P100 GPU는 각각 최대 2,048개의 액티브 쓰레드를 지원할 수 있는 56개의 스트리밍 프로세서를 보유하고 있다.

커널 코드가 전달된 어레이로 오프셋을 찾으려면, 해당 블록 및 쓰레드 인덱스를 알아야 한다. 병렬화된 커널은 다음과 같은 그리드 스트라이드(grid-stride) 루프를 사용하는 경우가 많다.

__global__
void add(int n, float *x, float *y)
{
int index = blockIdx.x * blockDim.x + threadIdx.x;
int stride = blockDim.x * gridDim.x;
for (int i = index; i < n; i += stride)
y[i] = x[i] + y[i];
}

CUDA 툴킷의 샘플을 보면 앞에서 다룬 기본 사항보다 더 많은 사항을 고려해야 한다는 것을 알 수 있다. 예를 들어 CUDA 함수 호출 중 일부는 checkCudaErrors() 호출을 통해 처리되어야 한다. 또한 대부분의 경우 가장 속도가 빠른 코드는 호스트 및 장치 메모리의 할당, 매트릭스의 복사와 함께 cuBLAS와 같은 라이브러리를 사용하게 된다.

요약하자면, 다양한 수준에서 GPU를 사용하여 애플리케이션을 가속화하는 것이 가능하다. CUDA 코드를 작성하고 CUDA 라이브러리를 호출할 수 있으며, 이미 CUDA를 지원하는 애플리케이션을 사용할 수도 있다.  editor@itworld.co.kr