#6 OpenCL 튜토리얼. Convolution Filter와 Local Memory 실습

__kernel
void convolution(__global float* imageIn,
    __global float* imageOut,
    __constant float* filter,
    int  rows,
    int  cols,
    int  filterWidth,
    __local float* localImage,
    int  localHeight,
    int  localWidth) {
 
 
    // 필터에서 사용하는 패딩의 양을 결정
    int filterRadius = (filterWidth / 2);
    int padding = filterRadius * 2;
 
    //워크그룹의 결과 영역 크기를 결정
    int groupStartCol = get_group_id(0)*get_local_size(0);
    int groupStartRow = get_group_id(1)*get_local_size(1);
 
    // 로컬 id받기 
    int localCol = get_local_id(0);
    int localRow = get_local_id(1);
 
    // 글로벌 ID, 
    int globalCol = groupStartCol + localCol;
    int globalRow = groupStartRow + localRow;
 
    // 로컬 메모리로 데이터 캐시 
        for (int i = localRow; i < localHeight; i +=
        get_local_size(1)) {
 
        int curRow = groupStartRow + i;
 
        for (int j = localCol; j < localWidth; j +=
            get_local_size(0)) {
 
            int curCol = groupStartCol + j;
 
            if (curRow < rows && curCol < cols) {
                localImage[i*localWidth + j] =
                    imageIn[curRow*cols + curCol];
            }
        }
    }
    barrier(CLK_LOCAL_MEM_FENCE);
 
    // 콘볼루션 수행
    if (globalRow < rows - padding && globalCol < cols - padding) {
 
        //필터반지름만큼 지나서 실행
        float sum = 0.0f;
        int filterIdx = 0;
        int offset = localRow*localWidth + localCol;
 
        sum += localImage[offset + 0] * filter[filterIdx++];
        sum += localImage[offset + 1] * filter[filterIdx++];
        sum += localImage[offset + 2] * filter[filterIdx++];
        sum += localImage[offset + 3] * filter[filterIdx++];
        sum += localImage[offset + 4] * filter[filterIdx++];
 
        offset += localWidth;
 
        sum += localImage[offset + 0] * filter[filterIdx++];
        sum += localImage[offset + 1] * filter[filterIdx++];
        sum += localImage[offset + 2] * filter[filterIdx++];
        sum += localImage[offset + 3] * filter[filterIdx++];
        sum += localImage[offset + 4] * filter[filterIdx++];
 
        offset += localWidth;
 
        sum += localImage[offset + 0] * filter[filterIdx++];
        sum += localImage[offset + 1] * filter[filterIdx++];
        sum += localImage[offset + 2] * filter[filterIdx++];
        sum += localImage[offset + 3] * filter[filterIdx++];
        sum += localImage[offset + 4] * filter[filterIdx++];
 
        offset += localWidth;
 
        sum += localImage[offset + 0] * filter[filterIdx++];
        sum += localImage[offset + 1] * filter[filterIdx++];
        sum += localImage[offset + 2] * filter[filterIdx++];
        sum += localImage[offset + 3] * filter[filterIdx++];
        sum += localImage[offset + 4] * filter[filterIdx++];
 
        offset += localWidth;
 
        sum += localImage[offset + 0] * filter[filterIdx++];
        sum += localImage[offset + 1] * filter[filterIdx++];
        sum += localImage[offset + 2] * filter[filterIdx++];
        sum += localImage[offset + 3] * filter[filterIdx++];
        sum += localImage[offset + 4] * filter[filterIdx++];
 
 
        // 데이터 쓰기 
        imageOut[(globalRow + filterRadius)*cols +
            (globalCol + filterRadius)] = sum;
    }
 
    return;
}

padding값 구하기

  5x5의 가우시안 필터를 사용하는 예제입니다. 이전 예제에서 본 것과 마찬가지로 Global Memory의 데이터를 Local memory로 옮길때 양쪽으로 2만큼씩의 padding을 더해 주어야 합니다. 그래야 Local Memory 내의 모든 픽셀이 주변 5x5크기의 픽셀들을 이용해서 값을 결정 할 수 있습니다. 

    // 필터에서 사용하는 패딩의 양을 결정

    int filterRadius = (filterWidth / 2);

    int padding = filterRadius * 2;

Global Memory와 Local Memory의 인덱싱 및 캐시

  이전 강좌에서 Global Memory 에서 인덱싱 하는 방법에 대해서 익혔습니다. 그렇다면 Global Memory와 Local Memory가 함께 있을때는 어떻게 메모리 할당을 해야 할까요? 간단히 그림으로 살펴보겠습니다. 

  빨간점에 해당하는 작은 사각형의 개수가 group_id입니다. 현재 (2, 1)위치에 있죠. 작은 사각형 하나가 Local Memory의 크기입니다. Local Memory의 크기를 group_id에 곱해준다면 실제 Global Memory에서 빨간 점의 위치를 구할 수 있습니다. 그 위치에 Local_id의 x 값과 y값을 각각 더해준다면 파란점의 위치. 다시말해 Global Memory내에서 연산의 대상이되는 픽셀의 위치를 구할 수 있습니다. 그림과 코드를 함께 보시면 금방 이해가 되실 꺼에요. 

    //워크그룹의 결과 영역 크기를 결정

    int groupStartCol = get_group_id(0)*get_local_size(0);

    int groupStartRow = get_group_id(1)*get_local_size(1);

 

    // 로컬 id받기 

    int localCol = get_local_id(0);

    int localRow = get_local_id(1);

 

    // 글로벌 ID, 

    int globalCol = groupStartCol + localCol;

    int globalRow = groupStartRow + localRow;

  가장 중요하다고 할 수 있는 로컬메모리로 데이터를 캐시하는 부분입니다. localHeight와 localWidth의 크기만큼 로컬 메모리에 집어넣습니다. imageIn은 입력받은 이미지, 곧 Global Memory입니다. Global Memory에서의 인덱스 curRow와 curCol과 Global Memory의 너비 cols를 이용해 픽셀의 위치를 구하고 이를 Local Image에 입력합니다. 

for (int i = localRow; i < localHeight; i +=

        get_local_size(1)) {

 

        int curRow = groupStartRow + i;

 

        for (int j = localCol; j < localWidth; j +=

            get_local_size(0)) {

 

            int curCol = groupStartCol + j;

 

            if (curRow < rows && curCol < cols) {

                localImage[i*localWidth + j] =

                    imageIn[curRow*cols + curCol];

            }

        }

    }

  그리고 로컬메모리의 각 픽셀에 대해 주변 픽셀에 Filter값을 곱해 픽셀의 값을 결정해줍니다. 

  루프를 풀어서 사용하는 루프언롤링 기법으로 작성한 코드입니다. 그냥 넘어가기가 애매하여 자세한 설명과 비슷한 다른 최적화 기법들과 함께 다음 포스트에서 다루면서 마무리 하도록 하겠습니다.