Gaussian Blur (OpenCL 사용, GPU 에서)

Gaussian Blur (OpenCL 사용!)

지난 포스팅 (CPU Gaussian Blur) 이후로 Gaussian Blur 을 OpenCl 로 GPU 에서 수행되도록 하는 실습을 진행할 것이다.

지난 포스팅은 아래 링크에서 볼 수 있습니다.

https://sh1mj1-log.tistory.com/21

또한 OpenCL 환경에 대해서도 알고 있어야 이번 실습이 가능하기 때문에 OpenCL 에 대해 모르신다면 아래 포스팅도 보고 오시면 도움이 될 겁니다!

https://sh1mj1-log.tistory.com/20

바로 실습에 들어가 봅시다.

실습

GPU Gaussian Blur

GaussianBlurOpenCL 동작 과정

동작 과정은 위와 같다.

Blur.cl 이라는 kernel 파일을 읽고, bmp 파일을 읽은 후, OpenCL 을 구동하기 위해 필요한 설정을 초기화한다.

그리고 kernel 의 수행을 위한 버퍼 등을 설정한 후 Launch kernel 한다.

kernel 실행 후 디바이스에서 수행된 결과가 저장된 buffer 을 Host 로 다시 가져온다.

이 코드를 정상적으로 실행하기 위해서는 아래의 것들을 먼저 준비해야 한다.

• ImageProcessing.h
• bmp.c
• lena.bmp
• GaussianBlurGPU.cpp
• Blur.cl
• Makefile

GaussianBlurGPU 구조

ImageProcessing.h 와 lena.bmp(이미지 비트맵 파일)은 CPU Gaussian Blur 에서와 완전히 똑같은 파일이다.

먼저 Blur.cl 파일을 봅시다!

Blur.cl

아직 OpenCL kernel 에 익숙하지 않으니 천천히 봅시다.

먼저 #pragma 는 precompile 되는 구문이다. OpenCL 에서 double 을 사용한다면 double precision 데이터 타입이 선언되기 전에 #pragma OPENCL EXTENSION cl_khr_fp64 : enable 을 넣어야 한다.

(참고)Double-precision floating-point format 은 더 정밀한 범위의 부동 소수점 포맷이다. 일반적으로 컴퓨터 메모리에서 64비트를 차지한다.

kernel_blur 함수에서 일단 row 에 1차원에서의 work-item의 ID 을 width 로 나눈 몫을 저장한다. 그리고 col 에 같은 work-item 의 ID 을 width 로 나눈 나머지를 저장한다. 이렇게 인수를 벡터로 선언해준다. row 와 col 을 설정하는 부분이 잘 이해가 되지 않는데 아래 그림과 함께 천천히 봅시다.

 

예를 들어서 width 가 4 pixel 이라고 합시다. 각 pixel 에 work-item 이 대응해서 Gaussian Blur 을 수행한다고 하면 아래처럼 동작하여 row 와 column 이 설정될 것이다.

물론 위처럼 영역이 완전히 구분되어 나누어지지 않고 저번 포스팅에서 설명했던 그림처럼 겹쳐서 수행될 것이다.

그리고 CPU Gaussian Blur 에서 했던 것처럼 float 형 2차원 배열을 만든 후 반복문을 통해 source 의 pixel 내에 있는 red, green, blue 값을 변경해준다. 마지막으로 변경된 red 값을 적용해준다.

 

GaussianBlurGPU.cpp

GaussianBlurGPU.cpp (1)

checkCL 이라는 이름으로 앞으로 오류가 발생했을 때 체크할 매크로를 정의한다.

openCL 을 사용하기 위해 설정하는 함수와 kernel 을 launch 하기 위한 함수를 미리 정의해준다.

opencl_infra_creation 메서드 (GaussianBlurGPU.cpp)

GaussianBlurGPU.cpp open_infra_creation 구현부

이 함수에서 OpenCL 을 사용하기 위한 OpenCL infra 와 kernel infra 을 만든다.

clGetPlatformIDs - 시스템에서 OpenCL 이 동작하는 플랫폼을 찾아서 플랫폼의 ID을 가져온다.

clGetDeviceIDs - 플랫폼에서 사용할 디바이스 정보를 불러온다.

clCreateContext - 디바이스로부터 OpenCL context 을 생성한다.

clCreateCommandQueue 특정 디바이스에 command-queue 을 생성한다.

clCreateProgramWithSource - OpenCL 커널 소스코드로 프로그램 객체를 생성한다. context의 program object을 생성하고 문자열 배열의 소스코드를 program object로 로드한다.

clBuildProgram - 프로그램 객체의 소스코드를 program executable 로 런타임 빌드한다.

clCreateKernel - 커널 오브젝트를 생성한다.

 

opencl_infra_creation 메서드 GaussianBlurGPU.cpp

GaussianBlurGPU.cpp launch_the_kernel 함수 구현부

clCreateBuffer - 커널이 동작할 때 처리할 데이터가 저장되는 (디바이스 버퍼)메모리 객체를 생성한다. source 버퍼와 결과가 저장될 버퍼를 생성한다.

clEnqueueWriteBuffer - Host 메모리에서 buffer object(디바이스 메모리)에 write 하기위해 enqueue 한다.

clSetKernelArg - 커널에서 처리할 데이터가 저장되어 있는 메모리 객체 정보를 커널에 추가한다.

clEnqueueNDRangeKernel - 디바이스에서 커널을 실행하기 위해 커맨드를 enqueue 한다.

clFinish - 큐에 있는 모든 객체들이 비워질 때까지 대기한다.

clEnqueueReadBuffer - 메모리에 저장되어 있는 데이터들을 읽어서 Host 배열에 저장한다.

main 함수 GaussianBlurGPU.cpp

GaussianBlurGPU.cpp main 함수

Blur.cl 이라는 커널 함수가 있는 소스파일을 읽는다. kernel 파일이 저장될 메모리 공간을 따로 할당하고 해당 메모리의 마지막 문자를 null문자(\0)로 설정한다.

blur 처리할 source 이미지인 lena.bmp 파일을 읽는다. blur 처리 결과가 저장될 메모리 공간을 따로 할당한다.

GaussianBlurGPU.cpp main 함수

각 local work group 에 있는 work-items 의 수(localSize)를 저장한다.

총 work-items 의 수(globalSize)도 결정한다. 이 때 실제 총 work-items 의 수가 local work-items 수의 정수배가 아닐 수 있으니 grid 부분에서 간단한 조건 연산을 수행한다.

그 후 우리가 구현해두었던 opencl_infra_creation 함수와 launch_the_kernel 함수를 수행한다.

저번 CPU Gaussian Blur 프로젝트와 똑같이 커널이 실행된 후의 결과를 write_bmp 함수를 통해 실제 결과 이미지에 적용을 해준다.

마지막으로 OpenCL 리소스와 host 의 memory 을 해제해준다.

Makefile

CC = arm-linux-androideabi-g++
ADB=adb

OPENCL_PATH =/code/W10/GaussianBlur_template/gpu
CFLAG =  -I$(OPENCL_PATH)/include -g
LDFLAGS = -l$(OPENCL_PATH)/lib/libGLES_mali.so -lm

TARGET=GaussianBlurGPU
TARGET_SRC =$(TARGET).cpp bmp.c

all: $(TARGET)

$(TARGET): $(TARGET_SRC)
	$(CC) -static $(TARGET_SRC) $(CFLAG) $(LDFLAGS) -fpermissive -o $(TARGET)
	echo
	echo "**** Install:" /data/local/tmp/$(TARGET)"****"
	$(ADB) push $(TARGET) /data/local/tmp
	$(ADB) push lena.bmp /data/local/tmp
	$(ADB) push Blur.cl /data/local/tmp
	$(ADB) shell chmod 755 /data/local/tmp/$(TARGET)

clean:
	rm -f *.o
	rm -f $(TARGET)

CFLAG 에서 gcc option 을 -I$(Directory_PATH) 로 주고 있다. (i대문자)

-I 옵션은 헤더 파일의 디렉토리를 추가하는 것이다.

또 -g 옵션은 GDB 디버거로 쓰이는 디버그 정보를 생성하는 것이다.

LDFLAGS 에서 gcc option 을 -l$(Directory_PATH) 로 주고 있다.

-l 은 라이브러리 파일을 링크하는 옵션이다.

필요한 헤더 파일과 라이브러리를 올바른 경로로 지정하여 컴파일 해줍니다.

 

make 후 실행하는 과정은 저번 포스팅 CPU Gaussian Blur 와 같다.

(HOST)
make
adb shell

(DEVICE)
LD_LIBRARY_PATH=/vendor/lib/egl
cl /data/local/tmp
./GaussianBlur

컨르롤 C

(HOST)
adb pull /data/local/tmp/blurred_lena.bmp ./
display blurred_lena.bmp

실습 결과

원래 이미지 lena.bmp

lena.bmp

 

디바이스로부터 호스트로 파일을 pull

아래는 blur 된 이미지! blurred_lena.bmp

blurred_lena.bmp

 

GPU Gaussain BLUR 에서 걸린 시간

위는 GPU Gaussian Blur 에서 걸린 시간이다!

그렇다면 저번 CPU Gaussian Blur 에서의 걸린 시간은 얼마나 들었을까??

CPU Gaussian Blur 에서 걸린 시간

바로 4028 정도 들었다. OpenCL 을 사용하여 GPU 로 처리를 한 결과 시간이 무려 18배 이상 차이가 났다!!

 

실습 GrayScale 과 Rotate

GPU 로 Gaussian Blur 을 수행했는데 당연히 다른 작업도 수행할 수 있다. OpenCL에서 디바이스의 GPU 을 이용하여 pixel 단위로 여러 반복 작업을 해야 하는 작업이면 무엇이든 가능하다.

기존 코드 구성에서 약간의 코드만 수정하여 GrayScale 을 적용하는 것과 시계 방향 90도 rotate 하는 실습도 진행해봅시다.

GrayScale

기존 Blur.cl 파일을 GrayScale.cl 파일로 대체한다.

GrayScale.cl

red, green, blue 값을 구글링 하여 어떠한 값으로 설정해야 할지 찾은 후 그 값으로 맞추어주기 위해 특정한 값을 곱하고 합하여 grayscale 의 색조로 만들었다.

 

또한 기존 GaussianBlurGPU.cpp 파일도 Grayscale.cpp 파일로 대체했다.

GrayScale.cpp

여기서는 커널로 사용될 커널 파일 이름과 커널 이름을 위에서 수정한 것과 같이 변경해준다. 그리고 결과로 얻을 이미지 bmp 파일의 이름을 graycale_lena.bmp 로 설정한다.

 

Makefile

당연히 Makefile 에서 경로와 이름도 바꾸어 주어야 한다~

CC = arm-linux-androideabi-g++
ADB=adb

OPENCL_PATH =/code/W10/GaussianBlur_template/grayscale
CFLAG =  -I$(OPENCL_PATH)/include -g
LDFLAGS = -l$(OPENCL_PATH)/lib/libGLES_mali.so -lm

TARGET=Grayscale
TARGET_SRC =$(TARGET).cpp bmp.c

all: $(TARGET)

$(TARGET): $(TARGET_SRC)
	$(CC) -static $(TARGET_SRC) $(CFLAG) $(LDFLAGS) -fpermissive -o $(TARGET)
	echo
	echo "**** Install:" /data/local/tmp/$(TARGET)"****"
	$(ADB) push $(TARGET) /data/local/tmp
	$(ADB) push lena.bmp /data/local/tmp
	$(ADB) push Grayscale.cl /data/local/tmp
	$(ADB) shell chmod 755 /data/local/tmp/$(TARGET)

clean:
	rm -f *.o
	rm -f $(TARGET)

Grayscale 결과

Rotate

똑같다. 기존 Blur.cl 파일을 rotate.cl 로 바꾸었다.

 

rotate.cl

여기서는 row/col/pix 와는 별개의 결과 drow/dcol/dpix 을 따로 만든다. d 는 dst(destination) 뜻일 겁니다~

시계 방향으로 90 도 돌리려면 row 와 col 을 바꾸면 된다.

// 기존에 <<행 곱하기 너비 + 열 >> 에서
pix = **row * width + col;**
		// << 열 곱하기 너비 + 행 >> 으로 변경해준다!
dpix = **dcol * width + drow**

 

똑같이 기존 GaussianBlurGPU.cpp 파일을 rotate.cpp 파일로 대체

rotate.cpp

이름을 바꿔준다!

 

Makefile

CC = arm-linux-androideabi-g++
ADB=adb

OPENCL_PATH =/code/W10/GaussianBlur_template/rotate
CFLAG =  -I$(OPENCL_PATH)/include -g
LDFLAGS = -l$(OPENCL_PATH)/lib/libGLES_mali.so -lm

TARGET=rotate
TARGET_SRC =$(TARGET).cpp bmp.c

all: $(TARGET)

$(TARGET): $(TARGET_SRC)
	$(CC) -static $(TARGET_SRC) $(CFLAG) $(LDFLAGS) -fpermissive -o $(TARGET)
	echo
	echo "**** Install:" /data/local/tmp/$(TARGET)"****"
	$(ADB) push $(TARGET) /data/local/tmp
	$(ADB) push lena.bmp /data/local/tmp
	$(ADB) push rotate.cl /data/local/tmp
	$(ADB) shell chmod 755 /data/local/tmp/$(TARGET)

clean:
	rm -f *.o
	rm -f $(TARGET)

Makefile 에서도 이름과 경로를 바꾸어준다.

 

rotate 결과

두번의 간단한 추가 실습으로 grayscale 을 OpenCL GPU로 수행하는데는 140 의 시간, rotate 을 OpenCL GPU 로 수행하는 데는 714 의 시간이 드는 것을 확인할 수 있었다.

 

아래는 cpu, gpu 을 이용한 gaussian blur, grayscale, rotate 가 모두 포함되어 있는 프로젝트 깃허브 입니다.

https://github.com/sh1mj1/GaussianBlur_OpenCL

GitHub - sh1mj1/GaussianBlur_OpenCL