cuda template project

cuda를 설치하고 나면 기본적으로 제공되는 예제들 중에 template에 관해서 살펴보겠습니다.

뭐 시작하기전에 메모리 종류라던가 커널이라던가 이것저것 알아야 할 것들이 많지만 일단 돌려보고 그런건 차차 하면서 알아 나가는게 좋지 않을까 합니다.

이 카테고리는 제가 공부하고 정리하는 의미가 있기 때문에 크게 순서에 얽매이지 않고 그냥 생각나는대로 쓰고 있음을 밝혀둡니다.;;

그럼 우선 한번 실행 시켜보도록 하죠.
Visual C++을 열고 C:\Program Files\NVIDIA Corporation\NVIDIA CUDA SDK\projects 폴더로 가면 예제들이 쭉 있습니다. 한 50개 좀 넘는군요..
이 중에서 template라는 폴더로 가서 솔루션을 엽니다.

그 다음 F5를 눌러 프로그램을 실행시켜 봅니다.
별 문제 없이 CUDA환경이 설정되었다면 아래와 같은 화면이 나올 겁니다.

시간은 머신마다 틀리게 나올거고 저렇게 나오면 정상적으로 돌아간 겁니다.

그럼 이제 저 화면이 어떻게 출력되었나를 template프로젝트의 소스를 보면서 알아보도록 하겠습니다.

우선 visual C++의 솔루션탐색기에서  src를 보면 3개의 파일이 존재함을 알 수 있습니다.
template.cu, template_gold.cpp, template_kernel.cu 이것들이 CUDA 프로그래밍을 하기 위한 기본적인 파일들인 것입니다.

이 파일들의 관계를 제대로 알아야지 CUDA프로그래밍을 할 수 있습니다.
NVIDIA에서는 쉽다고 말하면서 예제로 matrix연산 하는걸 메뉴얼에 올려놨는데 쉽기는...ㅡㅡ;
메뉴얼만 보면 쉽네? 하고 생각할 수 있겠지만 저같은 평민은 그 뒤에 감춰진 kernel이라던가 cu파일들 보면 어렵기만 하더군요.

우선적으로 template_gold.cpp를 보면 computeGold라는 함수에서 reference를 계산하는걸 볼 수 있습니다.

void
computeGold( float* reference, float* idata, const unsigned int len)
{
    const float f_len = static_cast<float>( len);
    for( unsigned int i = 0; i < len; ++i)
    {
        reference[i] = idata[i] * f_len;
    }
}

reference를 계산하는 목적은 GPU를 이용해서 계산한 값이 CPU와 이용한 값과 같은지를 비교하기 위해서 CPU를 이용해서 참조용 결과를 만드는 겁니다. 물론, 이건 예제이기에 저런 부분이 들어가는 것이고 GPU결과 검증을 위해서 필요한 부분이지만 실제로 프로그램을 작성할때는 저 부분은 빼주어도 cuda프로그래밍에 문제가 되지는 않습니다.

실제 CUDA프로그램은 template.cu, template_kernel.cu 라는 두 파일을 통해서 이루어 집니다. 이중에 template_kernel.cu는 실제 GPU상에서 연산이 이루어지는 내용을 포함합니다. 그리고 template.cu는 template_kernel.cu에 있는 testKernel이라는 함수를 돌리기 위한 준비과정과 돌리고 난후의 후처리를 담당합니다. 나중에 되면 굳이 파일을 2개로 분리할 필요없이 template.cu와 template_kernel.cu를 합쳐서 돌려도 별 문제없이 잘 돌아가는걸 확인할 수 있습니다. 일단, 개념 정립을 위해서는 template.cu는 GPU연산을 위한 준비와 후처리, kernel은 실제 GPU연산 수행 정도로 알아두면 될 것 같습니다.

template.cu를 보면 runTest라는 함수를 확인 할 수 있습니다.

void
runTest( int argc, char** argv)
{

    CUT_DEVICE_INIT();

    unsigned int timer = 0;
    CUT_SAFE_CALL( cutCreateTimer( &timer));
    CUT_SAFE_CALL( cutStartTimer( timer));

    unsigned int num_threads = 32;
    unsigned int mem_size = sizeof( float) * num_threads;

    // allocate host memory
    float* h_idata = (float*) malloc( mem_size);
    // initalize the memory
    for( unsigned int i = 0; i < num_threads; ++i)
    {
        h_idata[i] = (float) i;
    }

    // allocate device memory
    float* d_idata;
    CUDA_SAFE_CALL( cudaMalloc( (void**) &d_idata, mem_size));
    // copy host memory to device
    CUDA_SAFE_CALL( cudaMemcpy( d_idata, h_idata, mem_size,
                                cudaMemcpyHostToDevice) );

    // allocate device memory for result
    float* d_odata;
    CUDA_SAFE_CALL( cudaMalloc( (void**) &d_odata, mem_size));

    // setup execution parameters
    dim3  grid( 1, 1, 1);
    dim3  threads( num_threads, 1, 1);

    // execute the kernel
    testKernel<<< grid, threads, mem_size >>>( d_idata, d_odata);

    // check if kernel execution generated and error
    CUT_CHECK_ERROR("Kernel execution failed");

    // allocate mem for the result on host side
    float* h_odata = (float*) malloc( mem_size);
    // copy result from device to host
    CUDA_SAFE_CALL( cudaMemcpy( h_odata, d_odata, sizeof( float) * num_threads,
                                cudaMemcpyDeviceToHost) );

    CUT_SAFE_CALL( cutStopTimer( timer));
    printf( "Processing time: %f (ms)\n", cutGetTimerValue( timer));
    CUT_SAFE_CALL( cutDeleteTimer( timer));

    // compute reference solution
    float* reference = (float*) malloc( mem_size);
    computeGold( reference, h_idata, num_threads);

    // check result
    if( cutCheckCmdLineFlag( argc, (const char**) argv, "regression"))
    {
        // write file for regression test
        CUT_SAFE_CALL( cutWriteFilef( "./data/regression.dat",
                                      h_odata, num_threads, 0.0));
    }
    else
    {
        // custom output handling when no regression test running
        // in this case check if the result is equivalent to the expected soluion
        CUTBoolean res = cutComparef( reference, h_odata, num_threads);
        printf( "Test %s\n", (1 == res) ? "PASSED" : "FAILED");
    }

    // cleanup memory
    free( h_idata);
    free( h_odata);
    free( reference);
    CUDA_SAFE_CALL(cudaFree(d_idata));
    CUDA_SAFE_CALL(cudaFree(d_odata));
}

우선 CUT_DEVICE_INIT(); 를 호출하여 GPU를 사용할 준비를 합니다.
그 다음라인에서 수행시간을 측정하기 위한 변수를 할당하고 초기화 합니다.
     unsigned int timer = 0;
    CUT_SAFE_CALL( cutCreateTimer( &timer));
    CUT_SAFE_CALL( cutStartTimer( timer));

그리고 몇 개의 쓰레드를 돌릴 것인지 지정하고 거기에 맞는 메모리 크기를 지정합니다.
여기서 중요한게 쓰레드의 갯수를 몇개로 지정할 것인가인데 나중에 이부분이 계산 수행속도에 큰 영향을 미치게 됩니다.
    unsigned int num_threads = 32;
    unsigned int mem_size = sizeof( float) * num_threads;

그 다음 호스트와 디바이스에 각각 메모리를 할당하는 부분이 나오는데 호스트는 일반적인 CPU가 사용하는 ram을 말하는 것이고 디바이스 메모리할당은 그래픽카드에 있는 메모리를 사용하기 위해 할당하는 것입니다.
이때, cudaMalloc라는 함수를 이용해서 디바이스 메모리를 할당하고 cudaMemcpy라는 함수를 이용하여 호스트메모리에 있는 내용을 디바이스 메모리로 가져옵니다.
    // allocate host memory
    float* h_idata = (float*) malloc( mem_size);
    // initalize the memory
    for( unsigned int i = 0; i < num_threads; ++i)
    {
        h_idata[i] = (float) i;
    }

    // allocate device memory
    float* d_idata;
    CUDA_SAFE_CALL( cudaMalloc( (void**) &d_idata, mem_size));
    // copy host memory to device
    CUDA_SAFE_CALL( cudaMemcpy( d_idata, h_idata, mem_size,
                                cudaMemcpyHostToDevice) );

그리고 계산결과를 저장할 디바이스 메모리를 할당합니다.
    // allocate device memory for result
    float* d_odata;
    CUDA_SAFE_CALL( cudaMalloc( (void**) &d_odata, mem_size));

그 다음 CUDA에 존재하는 built-in 변수인 dim3형으로 그리드와 쓰레드를 지정해 줍니다. 이 부분은 nVidia 그래픽카드의 메모리구조를 알아야 이해할 수 있는 부분인데 그냥 아까 지정했던 갯수만큼 사용할 쓰레드를 지정한다고 생각하면 될 듯 합니다.(이 부분에 대한 정확한 내용은 아직 잘 모르겠습니다. 좀 더 공부 필요..근데 이런게 쉽다고하다니...메뉴얼 어디에도 설명이 없는데..-_-+)
    // setup execution parameters
    dim3  grid( 1, 1, 1);
    dim3  threads( num_threads, 1, 1);

그 다음 드디어 kernel을 실행시킵니다.
    // execute the kernel
    testKernel<<< grid, threads, mem_size >>>( d_idata, d_odata);

커널을 실행한 다음에 문제없이 잘 실행됐는지 확인을 합니다.
    // check if kernel execution generated and error
    CUT_CHECK_ERROR("Kernel execution failed");

호스트메모리에 디바이스에 저장되어 있는 결과를 가져올 메모리를 할당한 다음 디바이스에 있는 메모리의 결과정보를 호스트 메모리로 가져옵니다.
   // allocate mem for the result on host side
    float* h_odata = (float*) malloc( mem_size);
    // copy result from device to host
    CUDA_SAFE_CALL( cudaMemcpy( h_odata, d_odata, sizeof( float) * num_threads, cudaMemcpyDeviceToHost) );

그 다음 수행시간을 측정해서 화면에 출력합니다. 제일처음에 보았던 시간 출력이 이 부분입니다.
    CUT_SAFE_CALL( cutStopTimer( timer));
    printf( "Processing time: %f (ms)\n", cutGetTimerValue( timer));
    CUT_SAFE_CALL( cutDeleteTimer( timer));

reference용으로 쓸 결과를 CPU에서 계산 합니다.
    // compute reference solution
    float* reference = (float*) malloc( mem_size);
    computeGold( reference, h_idata, num_threads);

GPU에서 계산되어진 결과와 CPU에서 계산되어진 결과가 일치하는지 비교합니다.
    // check result
    if( cutCheckCmdLineFlag( argc, (const char**) argv, "regression"))
    {
        // write file for regression test
        CUT_SAFE_CALL( cutWriteFilef( "./data/regression.dat",
                                      h_odata, num_threads, 0.0));
    }
    else
    {
        // custom output handling when no regression test running
        // in this case check if the result is equivalent to the expected soluion
        CUTBoolean res = cutComparef( reference, h_odata, num_threads);
        printf( "Test %s\n", (1 == res) ? "PASSED" : "FAILED");
    }

할당되었던 메모리를 해제합니다.
    // cleanup memory
    free( h_idata);
    free( h_odata);
    free( reference);
    CUDA_SAFE_CALL(cudaFree(d_idata));
    CUDA_SAFE_CALL(cudaFree(d_odata));

이상이 기본적인 template.cu의 구조입니다.

쓰다보니 그냥 주석을 한글로 옮기는 정도밖에는 안되는군요..;;
글이 길어져서 template_kernel.cu는 다음 포스트로 넘어갑니다.