 |
| 圖一 |
Gaussian kernel 的大小由 sigma 決定, Host 端的程式根據 sigma 決定 kernel size, 根據 gaussian filter 的定義計算 mask 中的每個常數, 將 mask 儲存在 OpenCL 的 image object. kernel size 以及 mask 會給 OpenCL kernel 進行 convolution 時使用. 用以計算 Gaussian mask 的程式片段如下:
void
CreateBlurMask(float sigma, int * maskSizePointer)
{
int maskSize = (int)ceil(3.0f*sigma);
g_mask.resize((maskSize*2+1)*(maskSize*2+1));
float sum = 0.0f;
float* pmask = &g_mask[0];
pmask += maskSize*(maskSize*2+1) + maskSize;
for(int a = -maskSize; a < maskSize+1; a++)
{
for(int b = -maskSize; b < maskSize+1; b++)
{
float temp = exp(-((float)(a*a+b*b) / (2*sigma*sigma)));
sum += temp;
pmask[a+b*(maskSize*2+1)] = temp;
}
}
// Normalize the mask
for(int i = 0; i < (maskSize*2+1)*(maskSize*2+1); i++)
g_mask[i] = g_mask[i] / sum;
*maskSizePointer = maskSize;
}
根據我前一篇文章 OpenCL介紹 作修改, 在 host 端程式代碼的初始化, 大致上增加了以下的部分.
調用 CreateBlurMask, 並且給定 sigma = 3.0
CreateBlurMask(3.0, &g_maskSize);
接著, 建立一個 image object, 並且將 Gaussian mask 的內容由 host 端的 memory 複製到 device 端的 memory:
imgObject[2] = clCreateImage2D(GPU_context, CL_MEM_READ_ONLY, &mask_fmt, g_maskSize*2+1,
g_maskSize*2+1, 0, NULL, NULL);
err = clEnqueueWriteImage(commandQueue, imgObject[2], CL_TRUE, origin,
mask_region, 0, 0, &g_mask[0], 0, NULL, NULL);
if(err != CL_SUCCESS)
{
cleanUpOpenCL(GPU_context, commandQueue, program, kernel,
memoryObjects, numberOfMemoryObjects,
imgObject, numberOfImgObjects);
return 0;
}
設定好送進 kernel function 的參數
clSetKernelArg(kernel, 2, sizeof(cl_mem), &imgObject[2]); // source image
clSetKernelArg(kernel, 3, sizeof(cl_int), &g_maskSize); // mask Size
Host 端處理每個 frame 的流程如下:
(1) 產生一份原始影像 scaling down 成較低解析度後的結果.
(2) 讓GPU對 較低解析度的影像 進行 convolution.
(3) 讀回結果.
(4) CPU 將讀回的結果scaling up 成和原始影像一樣解析度.
(5) CPU 根據 ROI, 將scaling up後的影像和原始影像合成.
 |
| 圖二 |
因為要 composition, 故需要產生一份 composition 使用的 mask. 此mask 在 ROI 內的值為 0xff,其餘部分皆為 0.
camera >> Img1;
Mat mask(Img1.rows, Img1.cols, CV_8UC1);
mask.setTo(0);
Mat roi(mask, Rect(140,100,340,260));
roi = 0xff;
composition 的代碼如下, ocl_img 是從讀回讀回並 scaling up 的結果. Img0 則是原始影像.
ocl_img *= 0.3;
ocl_img.copyTo( Img0, mask );
接下來看 Kernel Function 的內容:
__constant sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE |
CLK_ADDRESS_CLAMP |
CLK_FILTER_NEAREST;
__kernel void laplacian( read_only image2d_t src_image,
write_only image2d_t dst_image,
read_only image2d_t mask,
__private int maskSize )
{
int x = get_global_id(0);
int y = get_global_id(1);
int2 pos = {get_global_id(0), get_global_id(1)};
// Collect neighbor values and multiply with Gaussian
float4 fSum = 0.0f;
float4 fMask;
float4 fColor;
uint4 uiColor;
for(int a = -maskSize; a < maskSize+1; a++)
{
for(int b = -maskSize; b < maskSize+1; b++)
{
int2 coord = (int2)((a+maskSize, b+maskSize));
uiColor = read_imageui(src_image, sampler, pos + (int2)(a,b));
fColor = convert_float4(uiColor);
fMask = read_imagef(mask, sampler, coord);
fSum += fColor * fMask.x;
}
}
uiColor = convert_uint4 (fSum);
write_imageui(dst_image, pos, uiColor);
}
Kernel 這邊的代碼 根據 mask size 進行 convolution, 再將結果寫到 output image object.
