cuda template project(2)

여기부터는 template_kernel.cu 에 관한 내용입니다. 첫번째 포스트쓰고 이래저래 딴짓하느라 그 사이 CUDA2.0 이 나왔네요.

template_kernel.cu는 실제 GPU연산을 담당하는 함수가 들어가 있습니다.

__global__ void
testKernel( float* g_idata, float* g_odata)
{
  // shared memory
  // the size is determined by the host application
  extern  __shared__  float sdata[];

  // access thread id
  const unsigned int tid = threadIdx.x;
  // access number of threads in this block
  const unsigned int num_threads = blockDim.x;

  // read in input data from global memory
  // use the bank checker macro to check for bank conflicts during host
  // emulation
  SDATA(tid) = g_idata[tid];
  __syncthreads();

  // perform some computations
  SDATA(tid) = (float) num_threads * SDATA( tid);
  __syncthreads();

  // write data to global memory
  g_odata[tid] = SDATA(tid);
}

 여기서 void라는 것은 다른 프로그래밍 언어를 보아왔다면 많이들 보셨을 반환값이 없다는 그 뜻입니다. 그런데 여기에 __global__라는 구문이 붙어 있는데 이건 그래픽카드에 있는 여러가지 메모리종류중에서 global 메모리를 사용하겠다는 뜻입니다. CUDA프로그래밍 중에서는 상황에 따라서 어떤 메모리를 사용하느냐가 속도에 상당히 큰 영향을 미칩니다. 우선은 __global__이 들어가는 부분이 어떤종류의 메모리를 사용한다를 선언한다는 정도만 알고 넘어갑니다.

  // shared memory
  // the size is determined by the host application
  extern  __shared__  float sdata[];
여기서도 마찬가지로 __shared__는 어떤 메모리를 사용할까를 지정하는 것입니다.

  // access thread id
  const unsigned int tid = threadIdx.x;
  // access number of threads in this block
  const unsigned int num_threads = blockDim.x;
스레드아이디와 블럭아이디를 얻어옵니다. 이건 CUDA 병렬화와 관계가 있는데 CUDA는 데이터를 병렬로 처리하는데서 계산속도의 큰 향상을 얻습니다. 그러기 위해서 물리적으로 여러개의 계산 유닛이 그래픽카드안에 구조를 가지고 자리를 잡게 됩니다. 이러한 구조도 나중에 살펴보겠지만 위에 나오는 구문은 그 구조중에서 어떤부분을 사용할지를 얻어온다고 생각하면 되겠습니다. 예를 들어 for문을 돌려서 루프를 계산할때 1부터 10까지의 계산을 순차적으로 한다고 하면 1번 계산하고나서 2번 계산하고 그 다음에 3번 계산하고 이런식으로 진행이 되지만 CUDA에서는 그 각각의 계산들이 한꺼번에 동시에 진행이 되기 때문에 계산속도가 10분의 1로 줄어들게 됩니다. 물론, 그렇게 계산을 수행하기 위해서는 for문 안에서 돌아가는 계산들이 각각 의존적이지 않고 독립적으로 돌아가야 정상적으로 계산을 수행하게 됩니다.


  // read in input data from global memory
  // use the bank checker macro to check for bank conflicts during host
  // emulation
  SDATA(tid) = g_idata[tid];
  __syncthreads();
GPU메모리에 있는 데이터를 읽어와서 bank를 체크한다는데 이부분은 아직 메뉴얼을 봐도 정확히 잘 모르겠네요. nvidia의 데이비드 아저씨가 강의한 파워포인트 자료를 봐도 이부분은 헷갈리는 부분입니다. 영어가 짧아서 듣기는 안되고 강의자료만 보려니 그렇네요..ㅡㅡ;


  // perform some computations
  SDATA(tid) = (float) num_threads * SDATA( tid);
  __syncthreads();
GPU에서 실제 계산을 수행합니다. 여기서   __syncthreads(); 이 부분이 있어야 전체 데이터들이 엉키지 않고 정상적으로 처리가 됩니다. 병렬화를 위해서 쪼개어져 계산되었던 값들이 모두 계산을 끝낼때까지 기다린다음에 값을 취합하는 겁니다. 만약 이 부분이 없다면 계산이 모두 끝나지 않았는데도 다음 구문으로 넘어가게 되니 계산 결과에 심각한 에러를 초래할 수 있습니다.

  // write data to global memory
  g_odata[tid] = SDATA(tid);
그 다음에 사용했던 데이터를 global 메모리로 옮깁니다. 이건 GPU상에 존재하는 여러가지 메모리중 global로 값을 옮겨서 계산이 끝난 결과값에 CPU에서 접근이 가능하도록 만드는 겁니다

이 정도면 대략적으로 CUDA에 포함되어 있는 template project의 소스는 살펴본 것 같습니다. 지루한 개념공부보다는 일단 뭐라도 코드를 한번 돌려보는게 어떨까 싶어서 이걸 제일 먼저 다뤘습니다. 다음부터는 CUDA 프로그래밍에 필요한 개념정리와 디버깅 방법에 대해 정리해 볼까 합니다.