#3 OpenCL 튜토리얼. 버퍼의생성과 커널의실행.

OpenCL 에서 커널이란 GPU에서 동작하는 함수 입니다. 그리고 커널은 GPU에 할당 된 메모리를 이용해 작업을 수행합니다. 

따라서 GPU에서 무언가 작업을 하기 위해서는 다음과 같은 작업 과정이 필요합니다. 

- 메모리 할당 및 데이터 쓰기

 GPU메모리를 할당하고 할당한 메모리에 값을 적재하는 부분 까지가 bufferWrite()함수의 내용입니다. 

 clCreateBuffer()함수로 이전 강의에서 생성한 context를 이용해 생성합니다. 인자로는 메모리를 Read용도로만 쓸 것인지, Write용도로만 쓸 것인지 둘다 사용할 것인지에 대한 인자와 사이즈 등이 들어있습니다. 버퍼를 내용을 변경하지 않는 constant버퍼로 사용할 경우 CL_MEM_READ_ONLY로 설정하면 읽는 속도가 더 빨라 진다고 합니다. 호스트가 메모리를 access할 수 있도록 할 수도 있고 기타 여러가지 기능이 있으니 다음 링크의 도큐먼트를 참고해 보세요.

https://www.khronos.org/registry/cl/sdk/1.0/docs/man/xhtml/clCreateBuffer.html

 

 버퍼를 생성 한 후 생성한 버퍼의 포인터와 넣을 값의 포인터를 인자로하는 clEnqueueWriteBuffer()함수를 이용해 GPU버퍼에 값을 써줍니다. 크기가 틀리게 되면 커널이 제대로 실행 되지 않는 오류가 생기니 나중에 큰 크기의 데이터를 처리할 떄는 꼭 주의해 주세요. 

 여기까지 완료하셨다면 이제 커널에서 CPU에서 업로드된 데이터를 사용할 준비가 다 된 것입니다. 

void bufferWrite()
{
    
 
    // GPU 메모리 버퍼 포인터 생성    
    d_inputArray_A = clCreateBuffer(context, CL_MEM_READ_WRITE,
        10*sizeof(int), NULL, NULL);
    d_inputArray_B = clCreateBuffer(context, CL_MEM_READ_WRITE,
        10*sizeof(int), NULL, NULL);
    d_outputArray = clCreateBuffer(context, CL_MEM_READ_WRITE,
        10 * sizeof(int), NULL, NULL);
    
    //GPU메모리에 넣을 배열값
    int inputArray_A[10] = { 2,2,2,2,2,3,3,3,3,3 };
    int inputArray_B[10] = { 1,1,1,1,1,4,4,4,4,4 };
 
    //배열값 출력 ㅜ분
    int i;
    printf("Array A : ");
    for (i = 0; i < 10; i++)
    {
        printf("%d ", inputArray_A[i]);
    }
    printf("\n");
    printf("Array B : ");
    for (i = 0; i < 10; i++)
    {
        printf("%d ", inputArray_B[i]);
    }
    printf("\n");
 
    //GPU 메모리에 데이터 
    clEnqueueWriteBuffer(queue, d_inputArray_A, CL_TRUE, 0, 10 * sizeof(int),
        inputArray_A, 0, NULL, NULL);
    clEnqueueWriteBuffer(queue, d_inputArray_B, CL_TRUE, 0, 10 * sizeof(int),
        inputArray_B, 0, NULL, NULL);
}

-OpenCL 커널 코드

 다음은 실제로 커널을 실행하는 부분입니다. 먼저 커널 코드부터 살펴보겠습니다. 단순히 두 배열을 더해서 같은 크기의 결과 배열에 넣는 코드입니다. 7행까지 뭔가 인덱스를 결정하고 결정한 인덱스를 결과에 집어넣는다. 정도만 인지하고 계신 채로 잠시 CPU에서 커널을 실행하는 코드로 넘어가겠습니다. 

__kernel
void simpleKernel(__global int* inputArray_A, __global int* inputArray_B, __global int* outputArray)
{
    uint dstYStride = get_global_size(0);
    uint globalRow = get_global_id(1);
    uint globalCol = get_global_id(0);
    uint dstIndex = globalRow * dstYStride + globalCol;
 
    outputArray[dstIndex] = inputArray_A[dstIndex]+inputArray_B[dstIndex];
}

 -커널을 실행하는 runKernel()함수

 6, 7줄 만 먼저 보도록 하겠습니다. 우리가 현재 사용하고 있는 배열은 1차원배열이며 10의 크기를 가지고 있습니다. 따라서 동시에 10개의 연산을 진행하기 위해서는 10개의 스레드를 생성해야 합니다. 

 단순히 10개의 스레드를 실행하는 것만으로는 계산을 할 수 없습니다. GPU커널에서는 현재 내가 몇번째 스레드인지를 알아야 합니다. 그래야 스레드 번호와 배열의 번호를 매칭시켜 배열의 모든 인덱스에 한번에 접근해 한번에 연산을 마칠 수 있게 됩니다. 

 globalSize[2]는 커널이 실행될 때 얼마나 큰 크기의 스레드를 한번에 실행 할 것인지를 2차원으로 나타내는 것이며, totalWorkItemsX와 totalWorkItemsY를 이용해 가로, 세로의 크기를 지정 해 줍니다.

void runKernel()
{    
    int i;
 
    //스레드 생성 개수 결정
    int totalWorkItemsX = 10;
    int totalWorkItemsY = 1;
    
    size_t globalSize[2] = { totalWorkItemsX, totalWorkItemsY };
 
    // 커널 매개변수 설정 
    clSetKernelArg(simpleKernel, 0, sizeof(cl_mem), &d_inputArray_A);
    clSetKernelArg(simpleKernel, 1, sizeof(cl_mem), &d_inputArray_B);
    clSetKernelArg(simpleKernel, 2, sizeof(cl_mem), &d_outputArray);
    
    //커널 실행
    clEnqueueNDRangeKernel(queue, simpleKernel, 2, NULL, globalSize,
        NULL, 0, NULL, NULL);
 
    // 완료 대기 
    clFinish(queue);
 
    //read 및 결과출력
    int outputArray[10];
    clEnqueueReadBuffer(queue, d_outputArray, CL_TRUE, 0,
        10*sizeof(int), outputArray, 0, NULL, NULL);
    
    printf("output  : ");
    for (i = 0; i < 10; i++)
    {
        printf("%d ", outputArray[i]);
    }
    printf("\n");
}

-커널내에서 스레드가 몇번째 위치인지 확인하는 방법

 커널 코드의 인덱스를 정하는 부분입니다. 위의 runKernel()함수 내의 globalSize가 GPU내에서 실행되는 스레드의 전체 크기 입니다. 

 get_global_size(0); 함수는 가로 크기를 반환합니다.

 get_global_id(1); 현재 스레드의 위치가 위에서부터 몇번째에 위치해 있는지 반환합니다. 

 get_global_id(0); 현재 스레드의 위치가 왼쪽에서부터 몇번째에 위치해 있는지 반환합니다. 

uint dstYStride = get_global_size(0);
uint globalRow = get_global_id(1);
uint globalCol = get_global_id(0);
uint dstIndex = globalRow * dstYStride + globalCol;

 따라서 위 정보들로 현재 스레드의 위치를 알 수 있고, 배열의 모양과 globalSize의 모양을 동일하게 설정 했으므로 현재의 스레드의 위치가 곧 배열의 인덱스가 됩니다. 

-커널 매개변수 설정과 커널의 실행

 커널의 매개변수 설정방법은 아래 그리과 같습니다. 커널에서 할당한 데이터에 맞는 형식의 데이터를 순서대로 매핑 해주시면 됩니다. 아까 생성한 메모리 버퍼의 포인터를 clSetKernelArg()함수를 이용해 매핑해주면 아까 배열에 넣었던 값을 커널에서 사용 가능합니다. clCreateBuffer를 통해 생성한 값 뿐만 아니라 일반 int, float 변수들도 변환이 불가능한 constant값으로 입력이 가능합니다. 예를들어 kernel code에서 내가 CPU에서 정의한 int값 만큼 결과값에 곱해주고 싶다면 

 clSetKernelArg(simpleKernel, 3, sizeof(int), &num); 과 같은 형태로 추가해주고, kernel함수의 인자에 4번쨰로 __global int*가 아닌 일반 int num등으로 인자를 추가해주시면 됩니다. 

 커널의 실행은 clEnqueueNDRangeKernel()을 통해 이루어집니다. 함수의 첫번째 인자인 queue는 이전 강의에서 생성한 queue입니다. 큐를 이용해서 커널을 실행하고 메모리를 매핑 하는등의 일을 합니다. 세번째 인자인 2는 스레드의 차원을 나타냅니다. 이 예제와 같이 1차원으로 된 예제는 세번쨰 인자를 1로 하고 globalSize를 x,y값을 갖는 변수가 아닌 하나의 값을 가지는 값으로 대체해서 사용이 가능 합니다. 다섯번쨰 인자인 localSize는 이후 있을 강의 중 자세하게 다루게 될 것입니다. 간단히 설명하면 global threads를 세분화 하여 각각 공통의 메모리 공간을 가지게 하고 빠르게 처리 할 수 있게 해주는 역활을 해준다고 생각하시면 됩니다. 

 clEnqueueNDRangeKernel(queue, simpleKernel, 2, NULL, globalSize,

        NULL, 0, NULL, NULL);

 clFinish()는 clFinish()이후의 코드가 실행 되기 전에 GPU에서 처리하도록 명령한 작업이 모두 끝났음을 보장 해주는 코드입니다. 만약 kernel의 실행이 완료 되지 않은 채 결과를 출력하려고 메모리의 값을 받아온다면 잘못된 값을 받아 올 수 도 있기 때문입니다. 

 clEnqueueWriteBuffer()를 통해 버퍼에 데이터를 쓴 것처럼 clEnqueueReadBuffer()를 통해 버퍼로부터 데이터를 읽어 올 수 있습니다. 

 두번의 강의에 걸쳐 기본적인 OpenCL 호스트 코드와 디바이스코드의 구조 및 사용방법을 알아 보았습니다. 앞으로는 재미있는 예제, gpu에서 활용 될 수 있는 다양한 기법들에 대해서 소개 해 보도록 하겠습니다.