CUDA Rolling을 이용한 Thread Optimization

- Rolling을 이용한 Thread 최적화 작업

CUDA로 연산을 병렬화로 구현하는 경우 GRID x THREAD 로 병렬화를 구분한다.

GRID 는 SM (Stream Multiprocessor)에 할당이 되고,
THREAD는 SM 안에 있는 SP (Stream Processor) 즉, CUDA Core에 할당된다.
SM안에 있는 SP 들은 공유 메모리 (Shared Memory)를 이용하여 빠른 속도로 서로 Data를 주고 받을 수 있다.

만약 2000 x 1000 의 작업을 병렬화로 구현해야 하는 경우 GRID : 2000, THREAD : 1000 개로 병렬화를 하면
가장 소스코드도 직관적으로 읽기 쉽고, 구현하기도 편하다.

SM 은 원래 서로 계산값의 영향을 미치지 않는 경우가 많아서 2000개로 나누어도 상관이 없다.
어차피 GPU는 SM 을 자신이 동시에 실행할 수 있는 갯수만큼 실행을 한 뒤 다음 SM들을 실행할 것이다.

하지만 THREAD의 경우는 Thread 간의 동기화 (Synchronize) 를 하는 Code도 들어가야 하며, 물리적으로 SM 안의 CUDA Core의 수가 넘어가는 Thread를 할당 할경우 Context Switching 도 발생하게 된다.
더군다나 GPU 가 한번에 실행시키는 Thread의 수는 생각보다 크지가 않다.

(GTX 770의 경우 동시 실행하는 SM은 8개, SM당 CUDA Core는 192개, WARP 는 32 이다.
 즉, SM 당 물리적으로 할당하는 Thread 수는 192 x 4 = 768개이며, 동시에 실행 가능한 Thread는 한번에 32개이다.)

그러므로 성능을 생각해서는 Thread의 수는 각 GPU의 WARP 만큼 하는것이 가장 이상적이다.
하지만, 굳이 성능을 위해서 머리아프게 Rolling을 해가면서 Code를 복잡하게 짜야할까 ?
굳이 성능이 아니더라도, 호환성을 위해서도 필요하다.
예를 들어서 위의 예제의 경우 Thread의 수를 1000개로 할당했는데, 다른 GPU의 경우 Thread 할당수가 200개 밖에 되지 않는 다면, 해당 Code로 실행을 하면 원하는 결과가 나오지 않는다.

아래의 Code를 참조하여 별로 어렵지 않게 Thread 수를 GPU의 WARP 만큼으로 설정을 하고 원래의 Code를 Rolling Code로 변환이 가능하다.
[-] Collapse

__global__ void Cuda_Test(float* i_fArray)
{
    int m = threadIdx.x - const_iNum[0] / 2;
    float fAng = (float)m * const_fPi[1] / const_iNum[1];
    i_fArray[threadIdx.x] = cosf(fAng) * cosf(fAng) + sinf(fAng) * sinf(fAng) - cosf(fAng);
}

[-] Collapse

__global__ void Cuda_Test(float* i_fArray)
{
    int idx = threadIdx.x;
    while (idx < const_iNumSensor[0])
    {
        int m = idx - const_iNum[0] / 2;
        float fAng = (float)m * const_fPi[1] / const_iNum[1];
        i_fArray[idx] = cosf(fAng) * cosf(fAng) + sinf(fAng) * sinf(fAng) - cosf(fAng);

        idx += blockDim.x;
    }
    __syncthreads();
}

과정을 말로 설명하자면 다음과 같다.

1. index를 하나 선언하여 기존의 threadIdx.x 의 자리에 넣는다.
2. while 문의 조건으로 index 가 기존 Thread의 크기보다 작을 경우 실행하도록 설정한다.
3. while 문 마지막에 index에 새로 설정한 Thread의 크기만큼 더한다.

4. while 문 뒤에 다른 연산이 더 있다면 __syncthreads(); 를 사용하여 Thread를 동기화 한다.