[Parallel Computing] (20) CUDA

CUDA는 아래의 정의와 같이 GPU를 General Purpose으로 사용할 수 있도록 하는 기술이다. 

CUDA("Compute Unified Device Architecture", 쿠다)는 그래픽 처리 장치(GPU)에서 수행하는 (병렬 처리) 알고리즘을 C 프로그래밍 언어를 비롯한 산업 표준 언어를 사용하여 작성할 수 있도록 하는 GPGPU 기술이다. CUDA는 엔비디아가 개발해오고 있으며 이 아키텍처를 사용하려면 엔비디아 GPU와 특별한 스트림 처리 드라이버가 필요하다. [1]

 

CUDA는 앞서 OpenCL과 같이 Host(CPU)와 Device(GPU)가 같이 연산을 수행하는 Heterogeneous Computing 환경에서 수행된다. OpenCL과 유사하게 Host program은 Device에서 수행할 Kernel의 실행을 관리한다. Kernel은 Compute device에서 수행되는 코드의 기본적인 단위이다. Host Program과 Kernel들은 병렬적으로 수행된다. 

Host(CPU)는 자기 자신의 메모리가 존재하고, Device도 마찬가지로 자기 자신의 메모리를 가지고 있다. Host Program에서 CPU 메모리에서 GPU 메모리로 데이터를 복사하고, GPU는 이 데이터를 로드하여 수행한다. 이때, GPU도 CPU와 같이 성능 향상을 위한 Cache를 가지고 있다. GPU에서 수행된 결과를 다시 CPU 메모리로 복사하여 CPU가 해당 결과를 사용한다. 

아래는 기초적인 CUDA 코드를 보여준다. 

"__global__" keyword는 해당 함수가 device에서 수행되는 것을 의미한다. 함수 Call은 host code에서 이뤄진다.

__global__ void mykernel(void) { }

int main(void) {
   mykernel<<<1,1>>>();
   printf("Hello World!\n");
   return 0;
}

-------------------------------------------------
- terminal
-------------------------------------------------
> nvcc hello_world.cu
> a.out
Hello World!
>

해당 코드는 'nvcc' Command로 컴파일을 수행하며, Device function은 NVIDIA Compiler가 처리하며, Host function은 Host compiler(eg, gcc)가 처리한다.

 

아래와 같이 두 개의 변수를 더하는 간단한 Kernel을 생각해보자.

__global__ void add(int *a, int *b, int *c){
 *c = *a + *b;
}

add() 함수를 device에서 수행하기 위해서는 아래와 같은 절차가 필요하다.

1. GPU 메모리에 Argument a, b, c를 저장할 공간을 할당해야 한다. 

2. Host 메모리에서 GPU 메모리로 데이터(a, b)를 전송한다. 

3. 연산을 수행한다.

4. 수행한 연산 결과(c)를 Host 메모리로 전송한다.

5. 메모리 할당을 해제한다.

CUDA API에는 Device 메모리를 관리하기 위한 함수들이 아래와 같이 존재한다.

>> cudaMalloc(), cudaFree(), cudaMemcpy()

 

CUDA는 SPMD(Single Program Multiple Data) 형태를 따르고 있다. 그러므로 위의 예제에서 add() kernel 하나의 프로그램으로 여러 데이터를 처리할 수 있어야 한다. OpenCL과 마찬가지로 각 Compute 단위마다 각자의 고유의 index를 만들 수 있다. CUDA는 가장 작은 단위인 Thread로부터 Thread의 모음인 Thread Block, Thread block의 모음인 grid로 이뤄져 있다. 아래와 같이 각 단위별로 크기와 idx를 얻을 수 있는 CUDA Built-in Variables들이 존재한다. 

// thread block id in the x/y/z-axis.
blockIdx.x/y/z

// thread id in the x/y/z-axis.
threadIdx.x/y/z

// thread block dimension(= # of threads in a block) in the x/y/z-axis
blockDim.x/y/z

// for example 
// 1d thread block = 4
// 1d thread per thread block = 8
index = threadIdx.x + blockIdx.x * 8(blockDim.x)

이를 통해서 각 Thread의 고유한 index를 부여할 수 있고, 이를 활용하여 아래와 같이 kernel를 작성할 수 있다. 

__global__ void add(int *a, int *b, int *c, int N){
  int idx = threadIdx.x + blockIdx.x * blockDim.x;
  
  // N = the size of a/b/c array.
  if(idx < N){
    c[idx] = a[idx] + b[idx];
  }
}

 

그러면 이 Kernel를 수행할 때, 사용하려고 하는 Thread와 Thread block은 어떻게 지정할 수 있을까? 아래 코드와 같이 function을 불러올 때, 지정할 수 있다. 첫 번째 항목은 # of Thread blocks, 두 번째 항목은 # of threads per a block이다. 

두번째 항목을 고정한고, 총 Thread의 개수가 연산하는 횟수인 N과 같도록 한다면 아래와 같이 표현 가능하다. 

// Launch add() kernel on GPU with N blocks
add<<<N/THREADS_PER_BLOCK, THREADS_PER_BLOCK>>>(d_a, d_b, d_c, N);

 

여기서 조금 디테일하게 고민해봐야 하는 점은 N과 THREAD_PER_BLOCK 간의 관계이다. Thread Block과 Thread는 정수이므로 N이 THREAD_PER_BLOCK으로 나누어 떨어지지 않으면, 반올림을 해야 한다. 이때, floor 연산이 이뤄질 경우, 전체 Thread 개수는 N보다 작아지는 경우가 발생하므로 모든 연산이 수행되지 않을 수 있다. 이를 방지하고자 아래와 같이 변경하면 N의 크기와 관계없이 모든 연산이 수행되는 것을 보장할 수 있다.

// N : the number of array elements
// M : blockDim.x
add<<<(N+M-1)/M, M>>>(d_a, d_b, d_c, N)

 

 

 

 

 

* Reference

[1]: https://ko.wikipedia.org/wiki/CUDA