CUDA소개
홈 > CUDA > CUDA소개

CUDA소개

집에서 또는 연구실에서 사용하고 있는 Desk Top 컴퓨터에 그래픽 카드 한 장을 추가함으로 인해서 계산 및 Simulation 속도가 200배가 빨라 진다면 과연 사람들은 이 기술을 어떻게 받아 들일까? 엔비디아는 CUDA(Compute Unified Device Architecture)라는 기술을 2006 10월 웹에서 공개하면서 본격적으로 개인용 슈퍼 컴퓨터의 시대를 열어가려고 하고 있다.

 

엔비디아의 데이빗 커크(David Kirk) 박사는 병렬 컴퓨팅 솔루션을 소개하는 자리에서 만약에 당신의 컴퓨터가 20~30% Performance 우위를 보인다 해도 비용적 측면을 생각해서 컴퓨터를 교체할지 말지에 대한 고민을 할 것이다. 그런데 100~200배까지 빨라진다면? 이러한 상황에서도 새로운 기술을 도입하지 않는다는 것은 단지 그 사람이 그 사실을 몰랐거나 게으른 것이다라고 기술했다.

 

오늘날의 GPU High Performance Computing(HPC)에 있어서 가장 적합한 솔루션이라고 평가되고 있으며 세상에서 가장 빠르게 확산이 되고 있는 기술 중에 하나이다. 특히 학술, 계산 분야에 있어서의 Heterogeneous Computing을 통해 시스템을 최적화함으로써 CPU Operating System, Task 처리와 같은 순차적인 업무를 위주로, GPU Massive Data를 처리하게 된다. GPU CPU대비 10배 가까운 메모리 인터페이스 Speed 240개의 Core에서 동시에 Data를 처리함으로써 최대 200배 이상까지도 계산 속도를 높일 수 있으며 또한 시스템을 최적화 시킬 수 있다.

 

 

1. NVIDIA CUDA H/W 아키텍쳐의 이해

 

NVIDIA GPU G80계열의 GPU부터 CUDA를 지원하게 된다. 특히, G200계열로 아키텍쳐가 진보하며 과학계산에서의 요구를 대폭 수용하였다. 현재의 NVIDIA GPU 240개의 코어를 가지고 있고, 8개의 코어가 묶여 30개의 멀티프로세서를 가지고 있다. NVIDIA Tesla 시스템은 그래픽 출력이 없는 계산전용 하드웨어로, 4GB의 메모리가 장착되어있다.

 

 

GT200계열 GPU의 내부구조

 

하나의 멀티프로세서는 8개의 SP(코어) 1개의 DP 유닛을 가지고 있다. 특히, 16384개의 32-bit 레지시터와 16KB의 공유메모리를 하나의 멀티프로세서가 공유하여 사용하게 된다. 또한, DP 유닛은 표준 더블프리시즌을 지원한다. 

 

 

 

 

NVIDIA Tesla  C1060의 자세한 스펙은 다음과 같다. 하나의 GPU에서 Single Precision 1Tflops 정도의 성능을 낸다. Double Precision의 경우 SP 8 + SFU 2, 10개의 Single Precision 유닛과 대비하여 하나의 DP 유닛만 존재하기 때문에 Single Precision 대비 약 1/10 정도의 성능을 나타낸다.

Processor

1 x Tesla T10

Number of cores

240

Core Clock

1.296 GHz

Floating Point Performance

933 GFlops Single Precision

78 GFlops Double Precision

On-board memory

4.0 GB

Memory bandwidth

102 GB/sec peak

Memory I/O

512-bit, 800MHz GDDR3

Form factor

Full ATX: 4.736” x 10.5”

Dual slot wide

System I/O

PCIe x16 Gen2

Typical power

160 W

Tesla C1060 스펙

 

 

2. CUDA 소프트웨어  환경

 

CUDA를 작동하기 위해 필요한 소프트웨어는 크게 3가지이다. 설치는 CUDA Driver를 설치하고, CUDA Toolkit을 설치한 후 CUDA SDK 예제를 설치하면 된다.

 

용도

비고

CUDA

Driver

CUDA지원 그래픽카드의 드라이버

NVIDIA에서는 통합 그래픽카드 드라이버를 제공하므로, 최신 그래픽카드 드라이버를 설치한다. Emulation설치를 할 경우 Driver의 설치가 필요 없다.

CUDA

Toolkit

CUDA 컴파일러와 라이브러리

현재  CUDA 2.2 버전이 나와있다. CUDA는 하위버전이 호환되므로, 최신버전의 CUDA toolkit을 다운받도록 한다.

CUDA

SDK

CUDA 코드 샘플 예제와 유용한 함수들

 

 

 

다음 그림은 CUDA Toolkit CUDA SDK를 그림으로 쉽게 표현한 것이다. 쉽게 이해한다면, 하드웨어를 제어할 수 있는 그래픽 드라이버가 CUDA Driver이고, 이를 OpenGL이나 DirectX로 제어하기 위해 복잡한 코딩이 필요하였는데, 계산(물리엔진, 시뮬레이션 등)을 위해서 CUDA 언어를 이용하면 쉽게 프로그래밍을 할 수 있도록 해주는 컴파일러 명령어라고 보면 된다.

 

3. CUDA 컴파일 환경

 

기존의 개발환경에서는 GPGPU를 위해서 C/C++에서 그래픽 하드웨어를 컨트롤 하기 위해 Cg, DirectX, OpenGL 등의 쉐이더 명령어를 사용해야 한다. 이러한 방법의 단점은 그래픽에 대한 API의 복잡성으로 인하여 쉽게 프로그래밍 하기 어렵다는 단점이 있다. CUDA는 이러한 단점을 해결하기 위한 방법인데, 이를 위해 좀더 쉬운 명령어 셋을 C언어에 추가하고, 추가된 명령어는 NVCC를 통해 GPU컨트롤이 가능한 Assemble 코드인 PTX Assemble 코드로 변환하여 사용할 수 있도록 한다.

 

nvidia CUDA

컴파일 워크플로우

 

이러한 과정은 복잡해 보일 수 있는데, C언어 기반의 기존 컴파일러와 CUDA 명령어를 컴파일 할 수 있는 NVCC 컴파일러 두 개가 각 시스템에 맞는 오브젝트 파일이 만들어지고, 최종 결과물은 하나의 실행파일이 만들어 진다. 여기서 cudafe CUDA Front End의 약자로 사용되었다.

 

 

 

CUDA C/C++에 확장된 명령어 셋이다. 표준 C/C++ 이외의 CUDA만의 Syntax가 존재하는데, 이는 기존 C Compiler에서는 작동하지 않고 NVCC에서만 인식하여 컴파일하게 된다. 명령어 확장 셋으로는 크게 4가지로 나뉘는데, Kernel 실행 시 사용되는 Triple Bracket, 함수 관리 선행어, 메모리 관리 선행어, CUDA 내장 구조체와 변수들이다.

 

 

4. 메모리 관리 선행어

 

CUDA 프로그래밍에서는 항상 CPU메모리의 데이터를 GPU 메모리로 옮겨 작업을 실행시키고, 그 결과를 받아와야 한다. (아래 그림 참조) 이러한 작업을 위해 메모리를 할당하고, 메모리를 복사하는 함수를 제공한다. 특히, CPU쪽 메모리와 GPU쪽 메모리에 대한 제어를 해주어야 한다. 또한, 그래픽카드의 온칩 메모리, 온보드 메모리 등에 대한 설정을 위해 다음과 같은 메모리 선행어를 사용한다.

__device__    global 메모리에 변수가 만들어 진다.

__const__     global 메모리에 변수가 만들어 진다. 단 읽기 전용으로cache가 작동된다.

__shared__   shared 메모리에 변수가 만들어진다.

메모리 관리 선행어

 

cudaMalloc, cudaMemcpy 함수를 통해 제어한다 

메모리 제어

 

 

5. 함수 관리 선행어

 

CUDA함수는 크게 3가지로 나뉜다. Device함수, Global 함수, Host함수이다. 일단, Host함수는 C표준 함수로, 파일의 입출력, 실행 결과를 비교 출력하는 등의 CPU에서 사용되는 작업의 함수를 선언할 때 사용된다. .c .cpp파일에서 __host__를 사용하지 않아도 기본 설정이 Host 함수이기 때문에 생략해도 되지만 .cu파일에서 NVCC 컴파일러에게 정확한 정보를 제공하기 위해 함수 정의 시 __host__를 사용하여 함수를 정의하는 것을 추천한다. Global 함수는 여러 개가 있을 수 있지만 하나의 GPU에서 1번씩 작동된다. , multiGPU가 작동될 때만 pthread, OpenMP 등의 컨트롤을 통해 여러 개를 실행할 수 있겠지만, 하나의 GPU에서는 Global 함수는 순차적으로 실행된다. , 하나의 GPU는 하나의 Kernel을 실행하는데, 그 커널을 실행할 때 사용되는 함수가 Global 함수이다.  __global__ 함수의 재사용 시 cuInit()을 작동시키면 Device Global Memory에 저장된 모든 변수가 초기화됨으로 조심하여야 한다. CUDA 개발자간에 Virtual GPU에 대한 논의가 있었는데, 물리적인 1개의 GPU를 여러 개의 GPU로 가상으로 인식하는 기술에 대한 논의가 있었지만, 아직 지원하고 있지 않다. __device__ 함수는 __global__ 의 함수에서 다양하게 사용되는 함수들을 정의할 때 사용된다. 

__global__   kernel을 실행시키는 함수를 정의할 때 사용

__device__   device영역에서 GPU 병렬처리 명령어들로 구성된 함수를 정의할 때 사용

__host__    C/C++기본 코드에서 사용되는 함수를 만들 때 사용(생략해도 됨)

함수 관리 선행어

 

다음은 __global__ 함수를 통해 간단한 벡터 합을 계산하는 프로그램을 나타내고 있다. For loop문을 CUDA 병렬 프로그램으로 처리된 것을 확인할 수 있다.

 

__global__  vectorAdd(const float * a, const float * b, float * c)
{  // Vector element index
    int nIndex = blockIdx.x * blockDim.x + threadIdx.x;
    c[nIndex] = a[nIndex] + b[nIndex];
}

 

 

 

6. Triple Bracket

 

Triple Bracket이 가장 CUDA 다운 확장 명령이라고 보면 된다. 사실  __XXXX__ C언어 확장은 이미 많은 경우 사용되고 있다. Windows 환경에서의 프로그래밍에서 동적 라이브러리를 만들고 사용할 때 dll expert import 시에도 __declspec(dllimport) 등과  같은 __XXX 형태의 선행명령어가 쓰이기 때문이다. Brook 명령어 셋에서는 A[] 배열 대신 A<>을 사용하여 병렬처리에 쓰이는 경우도 있지만, CUDA <<< A , B >>> 처럼 Triple Bracket을 사용하는 경우는 매우 독창적인 방법이라고 할 수 있다. 특히, 이러한 명령 Bracket C/C++에서는 인식하지 못하고 오직 NVCC에서만 인식한다.  <<>>의 역할은 __global__ 함수를 실행시킨 후 블록과 쓰레드로 나누어 병렬로 실행시키는 역할을 한다.

 

 

7. CUDA S/W 아키텍쳐

 

하나의 GPU 커널은 Triple Bracket에 의해 실행되는데, 30개의 멀티프로세서, 240개의 코어에 병렬적으로 Block Thread가 할당되게 된다. 이때, 하나의 GPU Grid와 대응되고, 하나의 멀티프로세서는 여러 개의 블록과 대응된다. 멀티프로세서 내의 8개 코어는 Block 내의 Thread와 대응된다.

H/W

S/W

CPU server

Host

GPU chip

Kernel, Grids

Multiprocessor

Blocks

SP (core)

Threads

PCI-e slot

Memcpy()

4GB memory

Global memory

Local Memory

Register

Register

Shared memory

Shared memory

하드웨어와 소프트웨어 환경 용어 정리

 

특히, Block 1차원 2차원, Thread 1차원, 2차원, 3차원으로 표현 가능하다. 이러한 성질로 인하여 고유 인덱스를 통해, 배열, 행렬 혹은 이미지, 동영상 등의 데이터를 접근하는데 편리하다.

 


Block Thread

 

또한, Block Thread Triple Bracket이 실행될 때 멀티프로세서의 개수에 의해 하드웨어적으로 자동 스케줄링이 이루어진다. 따라서 2개의 멀티프로세서가 있는 GPU에서 4개의 멀티프로세서가 있는 GPU로 업그레이드를 할 경우 코드의 변화 없이 2배의 성능향상을 기대할 수 있다.

 

 

Block H/W 스케쥴링

 

CUDA의 프로그래밍은 SIMD 방식으로 프로그래밍을 하는데, 데이터는 Thread, Block, 그리드에 따라 공유 가능한 데이터가 틀리다. 하나의 멀티프로세서가 독립적인 Shared Memory를 가지고 있기 때문에 Block간에는 서로 다른 데이트를 갖게 된다. 특히, 각각의 Thread 는 이웃한 동일 블록 내의 Thread 와 통신을 하기 위해서는 Shared Memory Global Memory를 거쳐 통신을 해야 한다. 행렬연산이나 기타 데이터 통신이 빈번한 프로그래밍을 CUDA로 적용할 때는 이러한 메모리 억세스 구조를 정확히 파악해서 프로그래밍을 해야 한다.

 

 

메모리 접속 구조

 

8. 결론

 

지금까지 CUDA 하드웨어와 소프트웨어 아키텍쳐의 기본적인 내용을 간단히 살펴보았다. NVIDIA 그래픽카드를 활용한 CUDA Heterogeneous HPC 다양한 분야에서 응용이 가능하고, 특히 SIMD 기반의 병렬프로그래밍에서 탁월한 성능 개선 효과를 얻을 있다. 해외에서는 CUDA 이용한 Heterogeneous HPC 대한 관심이 많고, 특히, Harvard, MIT, UIUC, 도쿄 공업대학 등에서 활발한 연구활동을 진행하고 있다. 연구소 뿐만 아니라, 전기전자, 기계설계, 자동차, 에너지, 의료, 금융 다양한 산업분야에서 다양한 기업들에서도 CUDA 이용한 생산성 향상에 관심을 갖고 있다. 국내에는 아직 CUDA 기술이 국내에 보급단계이다.