[CUDA] 2. Grid-stride loop

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

지난 정리에서 <<<,>>>로 thread의 사용을 표현했다. 해당 함수를 특정 스레드에서 실행시키는 것이다. 하지만, GPU의 진정한 꽃은 parallel programming을 실현하는 것이다. 기존 add함수에 병렬 스레드를 구성하는 것을 연습해보자.

 

Streaming Multiprocessors (SMs)라고 불리는 것은 GPU내에서 thread block을 직접 실행하는 주체가 된다. GPU 모델마다 다른 수의 SMs unit을 가지고 있고, SMs가 많을수록 더 좋은 성능의 병렬처리를 보인다고 한다. (아직은 단순한 이해)

 

kernel이 실행되면, blcok이 생성되고, 생성된 block은 하나의 SMs에 할당된다. blcok안에 존재하는 여러 thread들은 SM에 의해서 warp단위로 실행이되며, SM은 내부에서 이를 스케줄링한다. 1개의 warp은 32개의 thread로 고정이 되어있다. 처음부터 너무 많이 알면 재미없으니까 이쯤 넘어가자.

 

병렬 쓰레드들을 이루는 블럭들은 grid로 묶인다. CUDA에서 제공되는 grid 내의 block index (blockIdx.x), block내의 thread 수(blockDim.x), block내의 thread index (threadIdx.x)를 이용하여, 위처럼 특정 스레드의 인덱스를 자동으로 계산할 수 있다.

 

 

Grid-Stride Loop

__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];
}

당연하게도, 처리해야 할 계산의 수는 보통 스레드 수를 넘기 때문에, 모든 데이터를 한번에 처리할 수 없다. 그래서 grid 내의 index를 이용해서, 해당 스레드가 stride만큼 간격을 두고, 여러 개의 데이터를 처리한다. 위와 같은 stride를 통한 loop 구조를 grid-stride loop라고 한다. 

 

 

단순히 GPU위에서 덧셈을 하고, 더 나아가서 병렬 처리를 구현해봤다. 여기서 들어야 할 의문은 그래서 병렬처리로 수치적으로 얼마나 큰 향상을 주는지 알 수 있는가? 다음 글에서는 커널의 프로파일링에 대해 짚을 예정이다.