2. CUDA에서 데이터 처리 (3)

CUDA에서 데이터처리 (2)에서는 GPU에서 동작하는 커널 함수를 생성하고 호출하는 방법을 알아보았다. 이번 시간에는 병렬 프로그래밍의 꽃인 스레드의 개념에 대해 설명하겠다.

 

CUDA에서 데이터 병렬 처리

1. 입력과 출력에 사용할 데이터를 PC 메모리에 할당

2. 입력과 출력에 사용할 데이터를 그래픽 메모리에 할당

3. 처리하고자 하는 값을 PC 메모리에 입력

4. PC 메모리에 있는 입력 데이터를 그래픽 메모리로 복사

5. 데이터를 분할하여 GPU로 가져옴

6. 수천개 이상의 스레드를 생성하여 커널 함수로 병렬처리

7. 처리된 결과를 병합

8. PC 메모리에 결과를 전송

9. 그래픽 메모리를 해제

10. PC 메모리를 해제

 

이전시간에 CUDA에서의 데이터 병렬처리 구조에 대해서 알아보았는데 여기서 주목해야 할 것은 6번 수천개 이상의 스레드를 생성하여 커널 함수로 병렬처리를 한다는 부분이다. CPU 환경에서 생성하는 스레드는 생성 즉시 실행 코드가 프로세서에서 실행된다. 생성된 스레드의 개수와 CPU의 코어 수가 동일하면 최적의 효율을 나타내지만, 코어수보다 많아지만 시분할 스케줄로 스레드가 동작하게되어 효율이 떨어진다. CUDA에서 스레드는 그래픽 카드가 발휘할 수 있는 성능 이내의 개수에서는 즉시 생성되어 코어에서 동작하게 된다. 설정된 스레드 생성 개수가 그래픽 카드의 최대 성능을 낼 수 있는 스레드 개수를 넘어서면, 효율이 떨어지지 않는 스레드 개수 만큼 실행하고 나머지는 대기하게 된다.

 

Kernel<<<block, thread>>>(a, b, c);

이전 게시글에 CUDA에서 GPU 메모리, 즉 device에서 작동하는 함수인 커널 함수를 host에서 호출하는 방법을 설명하였다. 이때 블록과 스레드를 사용하여 CUDA의 꽃인 병렬 프로그래밍을 할 수 있다. 

커널 함수를 호출할때 병렬 프로그래밍을 위해 블록과 블록당 스레드 수를 설정한다고 하였다. 이때 정해진 스레드 수를 바탕으로 커널 함수에서 threadIdx.x와 blockIdx.x로 해당 스레드의 index에 접근을 할 수있다. 

__global__ void add( int *a, int *b, int *c ) {
int index = threadIdx.x + blockIdx.x * blockDim.x;
c[index] = a[index] + b[index];
}

이를 사용해서 device에서 병렬로 덧셈을 하는 커널 함수 이다. 해당 코드를 보면 스레드의 index를 얻기 위해 blockIdx.x에 blockDim.x를 곱해준다. blockDim.x는 블록당 스레드의 수이며, 블록당 스레드의 수가 8일때 21번째 스레드의 index에 접근하고 싶다면 아래와 같은 방법으로 접근을 할 수 있다. 

M을 blockDim.x 라고 보면 된다.