First Kernel
이미지 convolution 과정
#include<iostream>
using namespace std;
void convolve(
float *pInput,
float* pFilter,
float* pOutput,
const int nInWidth, //input image height
const int nWidth, //output image width
const int nHeight, // output image height
const int nFilterWidth, //filter size
const int nNumThreads)
{
for (int yOut = 0; yOut < nHeight; yOut++)
{
const int yInTopLeft = yOut;
for (int xOut = 0; xOut < nWidth; xOut++)
{
const int xInTopLeft = xOut;
float sum = 0;
for (int r = 0; r < nFilterWidth; r++)
{
const int idxFtmp = r*nFilterWidth;
const int yIn = yInTopLeft + r;
const int idxIntmp = yIn*nInWidth + xInTopLeft;
for (int c = 0; c < nFilterWidth; c++)
{
const int idxF = idxFtmp + c;
const int idxIn = idxIntmp + c;
sum = sum + pFilter[idxF] * pInput[idxIn];
}
}
const int idxOut = yOut*nWidth + xOut;
pOutput[idxOut] = sum;
}
}
}
int main()
{
//..
}
OpenCL Kernel 적용
Image Convolution Using OpenCL™ – A Step-by-Step Tutorial Step 2
xOut과 yOut은 좌표에 해당한다.
위 코드와는 달리 이번 커널에서는 xOut과 yOut을
get_global_id()로 초기화 한다.
__kernel void Convolve(const __global float *pInput, __constant float * pFilter,
__global float *pOutput,
const int nInWidth,
const int nFilterWidth)
{
const int nWidth = get_global_size(0);
const int xOut=get_global_id(0);
const int yOut=get_global_id(1);
const int xInTopLeft=xOut;
const int yInTopLeft=yOut;
float sum=0;
for(int r=0; r<nFilterWidth;r++)
{
const int idxFtmp=r*nFilterWidth;
const int yIn=yInTopLeft+r;
const int idxIntmp=yIn*nInWidth+xInTopLeft;
for(int c=0;c<nFilterWidth;c++)
{
const int idxF=idxFtmp+x;
const int idxIn=idxIntmp+x;
sum = sum+ pFilter[idxF]*pInput[idxIn];
}
}
const int idxOut=yOut*nWidth+xOut;
pOutput[idxOut]=sum;
}
2. 현재 사용할 수 있는 디바이스들을 요청한다.
size_t listSize;
/* 처음으로 디바이스 list의 size를 얻는다 */
clGetContextInfo(context,CL_CONTEXT_DEVICES,..., &listSize);
/* 디바이스 리스트를 할당한다 */
cl_device_id devices = (cl_device_id*)malloc(listSize);
/* 디바이스 리스트의 데이터를 얻는다*/
clGetContextInfo(context, CL_CONTEXT_DEVICES,listSize,devices,...);
3. command queue를 만든다. kernel을 특정 디바이스에서 실행 시키기 위해서는 command queue에 집어 넣으면
순서대로 kernel를 실행한다.
cl_command_queue queue = clCreateCommandQueue(context, devices[0], ...);
4. cl kernel 파일을 읽는다.
cl_program program = clCreateProgramWithSource(context,1,&source,...);
5 Build program.
clBuildProgram(program, 1, devices, ...);
6. 커널 명시
cl_kernel kernel = clCreateKernel(program, "Convolve", ...);
Initialize OpenCL Buffers
3가지의 버퍼가 필요하다.
1). Input image
2). output image
3). convolution filter
1. cl_mem inputCL = clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_USE_HOST_PTR,
host-buffer-size, host-buffer-ptr,....);
(만약 디바이스가 gpu라면 input image buffer에 data를 복사한다.)
clEnqueueWriteBuffer(queue,inputCL, ..., host-buffer-ptr,....);
Excute OpenCL Kernel
커널 호출하는 대신에, cl 디바이스에서 커널을 실행한다.
1. 아큐먼트를 설정한다. 값들도 설정.
/*input buffer , arg 0*/
clSetKernelArg(kernel, 0, sizeof(cl_mem), (void*) &inputCL);
/*filter buffer, arg 1*/
clSetKernelArg(kernel, 1, sizeof(cl_mem),(void*)&filterCL);
/*output buffer, arg 2 */
clSetKernelArg(kernel,2,sizeof(cl_mem),(void*)&outputCL);
/*input image width, arg 3*/
clSetKernelArg(kernel, 3, sizeof(int), (void*) &nInWidth);
/*filter width, arg 4*/
clSetKernelArg(kernel, 4, sizeof(int), (void*) &nFilterWidth);
2. command queue에 커널을 enqueue 한다. kernel은 디바이스에서 실행될 것이고,
결과 값은 output buffer에서 이용할 수 있다.
data buffer 치수 안에서 pass를 해야한다.
그리고 계산하기 위한 아이템 수를 알아야한다. ( 픽셀 수인듯..)
clEnqueueNDRangeKernel(queue, kernel, data-dimensionality,..., total-work-size, work-group-size, ...);
Release OpenCL Buffers
clReleaseBuffer(inputCL);
Shutdown OpenCL
program 종료 전에 opencl resource 를 release 해야 한다.
clReleaseKernel(kernel);
clReleaseProgram(program);
clReleaseCommandQueue(queue);
clReleaseContext(context);