GPGPU - 2부

GPGPU란?? - 2부

GPGPU를 위한 GPU 구조

지난 1부에서 언급한 것 과 같이 GPU는 그래픽 처리를 위한 하드웨어이고 그래픽 처리는 대량의 data-level-parallelism 을 가진다.
따라서, 기본적으로 SIMD 형태의 구조를 가진다.
(SIMD = Single Instruction Multiple Data)

위 그림 처럼 하나의 instruction을 여러개의 ALU가 동시에 여러 데이터를 처리 하는 것을 SIMD 라고 하고, 이처럼 생겼다.

뭐 대충 이러하다. 사실 이 그림은 2008년에 나온 NVIDIA Tesla 아키텍처를 그린 그림인데, 요즘 나오는 GTX980의 경우는 대충 이거의 20배 정도 크다고 생각하면 된다.
CUDA core 라고 되어 있는 연두색 네모 하나가 ALU 하나라고 보면 되고, 8개의 코어가 16개 있으니 128개의 데이터가 한번에 처리될 수 있다.
자세한건 NVIDIA White paper에...

실제 동작

아주 간단한 예제 코드를 통해 살펴보자.

__global__ void vecAdd(double *a, double *b, double *c, int n){
int id = blockIdx.x*blockDim.x+threadIdx.x;
    c[id] = a[id] + b[id];
}

위 코드는 아주 간단한 CUDA 코드로 하나의 커널이다.
커널은 GPU 에서 실행되는 단위로 C언어의 함수처럼 생겼다.
id 는 동시에 동시에 도는 애들 중에서 어떤애 인지를 지칭한다. (저 맨 위 그림에서 vec_r1 의 id는 1이라고 보면 된다.)
이 코드는 a[i] + b[i] 를 모든 CUDA 코어에서 동시에 수행한다.
두 번째 GPU를 기준으로 보면 128개의 element를 한번에 할 수 있다. 이론적으로는...

CPU 에서의 동작과 비교를 해보자.
이 CUDA 코드를 CPU 에서 돌린다고 하면, 보통 for 문을 통해서 돌리게 되는데

for(i = 0; i < N; i++){
    c[i] = a[i] + b[i];
}

이런식으로 간단하게 작성할 수 있다.
CPU 에서 동작을 하면 GPU 에서 128개를 한번에 돌릴 수 있는 것을 128번이나 반복해서 해야한다.
(사실 CPU에서 SIMD unit 이 내장되어 있고, 그게 아니더라도 ALU가 동시에 여러개 수행 가능하기는 하다...)

뭐 어쨋든, 최근 GPU의 경우는 2000개 이상 코어를 가지고 있기 때문에 이런 data-level parallel 한 work에 대해서는 CPU보다는 훨씬 빠르다.