후로링의 프로그래밍 이야기

#9 OpenCL 튜토리얼. 히스토그램 본문

OpenCL

#9 OpenCL 튜토리얼. 히스토그램

후로링 2016. 12. 27. 19:00

히스토그램(Histogram)


  도수분포표를 그래프로 나타낸 것을 히스토그램이라고 합니다. 간단히 얘기하면 히스토그램은 해당 항목에 대한 빈도수를 구해 그래프로 나타낸 것이라고 할 수 있습니다. 히스토그램은 자료를 분석하는데 있어서 아주 기본적인 도구입니다. 



히스토그램 연산


  예를들어 1부터 10 사이의 임의의 정수로 이루어진 100의 크기를 가진 배열이 있다고 하면 histogram[10]의 배열에 각 정수의 빈도수를 저장 하는 것입니다. 간단하지만 대상 자료의 크기가 커진다면? 이미지의 해상도에 따라 적게는 100만부터 1000만번씩 연산을 하게 됩니다. 여러개의 데이터에 대해 하나의 연산을 적용하는것이기 때문에 OpenCL로 최적화 하기 아주 좋은 케이스 입니다. 



메모리 충돌


  스레드 프로그래밍을 배울때 가장 중요한 개념이 바로 Memory lock입니다. 여러개의 스레드가 동시에 메모리에 접근해서 값을 수정한다면 제대로된 결과를 기대하기 힘들게 되지요. GPU연산도 마찬가지입니다. GPU의 각 스레드는 순서상관없이 동시에 실행 되기 때문에 히스토그램 연산을 하나의 결과배열에 대해 적용한다면 제대로된 결과 값을 얻을 수 없습니다. 



Atomic 연산


  원자적 연산은 하나의 프로세스가 제대로 완료 되기까지 interruption되지 않는 연산을 말합니다. OpenCL에도 메모리 공간이 연산 중에 침범당하지 않도록 하는 연산이 있습니다. 아래와 같은 형태로 쉽게 사용 할 수 있습니다. volatile 키워드는 실행중 값이 언제든 바뀔 수 있으니 새로 참조하라는 뜻입니다. 


1
2
3
4
5
int atomic_add (volatile __global int *p, int val)
int atomic_sub (volatile __global int *p, int val)
int atomic_xchg (volatile __global int *p, int val)
int atomic_inc (volatile __global int *p)
int atomic_dec (volatile __global int *p)
cs



Atomic 연산의 문제점


  자, 위에서 소개한 atomic연산을 사용하면 결과값은 제대로 나오게 됩니다. 하지만 아주 커다란 문제가 생기게 됩니다. atomic을 이용해 작성한 프로그램은 더이상 병렬 프로그램이 아니게 됩니다. GPU 메모리에 순차적으로 접근해야 하기 때문이죠. CPU보다 성능이 낮은 GPU코어에서의 직렬연산은 CPU에서 실행하는것보다 한참 느립니다. 따라서 메모리 충돌을 일으키지 않는 메모리 접근 방법이 필요하게 됩니다.



메모리 충돌을 회피하는 방법


  각 로컬 메모리에 히스토그램 배열을 생성하고 그곳에서 atomic연산을 한다면 로컬 메모리로 얼마나 분할 되는지에 따라서 병렬도가 달라지게 되겠지요? 한 로컬 메모리당 얼마만큼의 전역 메모리를 계산 하게 할 수 있느냐가 중요합니다. 한 로컬 메모리에서 연산하는 전역 메모리의 값이 적을수록 계산은 병렬화 되겠지만 이후에 다시 히스토그램을 합치는 리덕션 과정을 거쳐야 하기 때문에 적절한 수치를 찾는 것이 중요하겠습니다. 


atomic_inc()함수의 인자로 전역 메모리를 넣으면 전역 원자 연산을 하지만, 로컬 메모리를 넣으면 지역 원자 연산을 수행하게 됩니다. 다시말해 스레드가 모든 영상에 접근하지만 실제 값의 저장은 local 메모리에 하게 되는 것입니다.  이후 barrier()를 통해 로컬 메모리에 fence를 치고 로컬메모리에 히스토그램을 저장하는 과정이 모두 끝나기까지 기다린 후 히스토그램을 합쳐주는 작업을 하게 됩니다. 


__kernel void histogram(__global struct gray * input, __global int * global_hist){
  int group_id = get_group_id(0);
  int num_groups = get_num_groups(0);
  int local_id = get_local_id(0);
  int local_size = get_local_size(0);

  volatile __local int histogram[256];

  int i;
  for(i=local_id; i<256; i+=local_size){
    histogram[i] = 0;
  }

  int rowNum, colNum, value, global_hist_offset

  for(rowNum = group_id; rowNum < 1080; rowNum+=num_groups){
    for(colNum = local_id; colNum < 1920; colNum += local_size){
      value = input[rowNum*1920 + colNum].i;
      atomic_inc(histogram[input]);
    }
  }

  barrier(CLK_LOCAL_MEM_FENCE);
  global_hist_offset = group_id * 256;
  for(i=local_id; i<256; i+=local_size){
    global_hist[global_hist_offset + i] = histogram[i];
  }

}


Comments