[App 개발] OpenCL 프로그래밍 예제
본문
아직 제 지식과 기술이 일천하여 완전히 최고 속도를 내지는 못하지만, 그래도 공부에는 도움이 될만한 소스가 아닐까 싶습니다.
Bilinear Interpolation 기법을 이용한 이미지의 확대/축소 및 회전 루틴입니다. 사실 이 내용은 코어 그래픽스나 코어 이미지의 affine transform 을 이용하면 훨씬 더 깔끔하게 처리할 수 있습니다. 하지만, 가끔 이렇게 간단한 루틴이 핸디하게 필요할 때도 있지요.
보간법에 대한 내용은 각자 래퍼런스를 찾아서 공부하시면 될 것 같고, CUDA 나 OpenCL 공부에 관심 있으신 분들의 경우, 일단은 대략 매트릭스 조작 예제에 대해서는 거의 한 커널 프로세스 당 한 셀 혹은 한 픽셀 단위로 작업을 나누어 주는 것이 거의 정석처럼 되어 있습니다. 사실 그 방법이 가장 효율이 높은 방법일 것이라 생각합니다.
아마도 이것은 그래픽 프로세서의 특성 및 가용 리소스 등에서 그 원인을 찾아볼 수 있지 않을까 생각되는데요, 제가 사용하고 있는 맥북프로의 엔비디아 9600M 은 어쩐 일인지 매뉴얼대로 잘 작동되지 않는 듯한 느낌을 받고 있습니다. 실제 애플 개발자 예제도 에러가 나는 것을 보면 제 실수는 아닌 것이 거의 확실해 보이고요.
그리고, 9600M 이 휴대용 제품에 해당하는 프로세서라서 그런 것일까요? 단밀도 실수 연산은 문제가 없는데, 배밀도 실수 연산에 대응하는 수학 함수 라이브러리가 지원되지 않는 것 같습니다. 당장 OpenCL 을 실험해볼 기계가 맥북 프로밖에 없어서 확신하기 어려운데요, 얼렁 돈 모아서 맥 프로라도 한 대 장만해서 결판을 내어야겠습니다. 집에서 쓰는 아이맥은 아예 OpenCL GPU 를 지원하지 않는군요.
어쨌든, 커널 소스를 올립니다. 뭐 별반 대단한 것 없으니 적당히 훑어 보시고, 대충 이렇게 발로 짜면 되는구나 하는 것만 구경해 보시면 되지 싶고요, 이 정도만 해도 대략 CPU 연산 대비 50 배의 이득을 얻을 수 있었다는 것으로도 충분히 만족스럽습니다. ^^
__kernel void bilinear_interpolation_ind(
__global float *source,
__const int width,
__const int height,
__const float angle,
__global float *target,
__const int tw,
__const int th )
{
int xDest = get_global_id( 0 ) ;
int yDest = get_global_id( 1 ) ;
float fraction_x, fraction_y, one_minus_x, one_minus_y ;
if( xDest < tw && yDest < th ) {
int x = xDest - tw / 2 ;
int y = th / 2 - yDest ;
float distance = sqrt( (float)( x * x + y * y ) ) ;
float polarangle = 0.0f ;
if( x == 0 ) {
if( y == 0 ) {
} else if( y < 0 ) {
polarangle = 1.5f * 3.14159265f ;
} else {
polarangle = 0.5f * 3.14159265f ;
}
} else {
polarangle = atan2( (float)y, (float)x ) ;
}
polarangle -= angle ;
float truex = distance * cos( polarangle ) ;
float truey = distance * sin( polarangle ) ;
truex = truex + (float)( tw / 2 ) ;
truey = (float)( th / 2 ) - truey ;
int floor_x = floor( (float)( truex * (float) width / (float) tw ) ) ;
int floor_y = floor( (float)( truey * (float) height / (float) th ) ) ;
int ceil_x = floor_x + 1 ;
if( ceil_x >= width ) ceil_x = floor_x ;
int ceil_y = floor_y + 1 ;
if( ceil_y >= height ) ceil_y = floor_y ;
if( floor_x < 0 || ceil_x < 0 || floor_x >= width || ceil_x >= width ||
floor_y < 0 || ceil_y < 0 || floor_y >= height || ceil_y >= height ) {
} else {
fraction_x = truex * (float) width / (float) tw - (float)floor_x ;
fraction_y = truey * (float) height / (float) th - (float)floor_y ;
one_minus_x = 1.0 - fraction_x ;
one_minus_y = 1.0 - fraction_y ;
float topleft = source[ floor_y * width + floor_x ] ;
float topright = source[ floor_y * width + ceil_x ] ;
float bottomleft = source[ ceil_y * width + floor_x ] ;
float bottomright = source[ ceil_y * width + ceil_x ] ;
float topvalue = one_minus_x * topleft + fraction_x * topright ;
float bottomvalue = one_minus_x * bottomleft + fraction_x * bottomright ;
float newvalue = one_minus_y * topvalue + fraction_y * bottomvalue ;
if( newvalue < 0.0 ) newvalue = 0.0 ;
if( newvalue > 1.0 ) newvalue = 1.0 ;
target[ yDest * tw + xDest ] = newvalue ;
}
}
}
최신글이 없습니다.
최신글이 없습니다.
댓글목록 0