圖一 |
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.
VICI, Casino, Slots, Table Games, - FilmfileEurope
回覆刪除VICI, 동인지사이트 Casino, Slots, Table 토토 사이트 샤오 미 Games, 우리형티비 Casino Games, Casino, Games, Casino, Games, Casino, 사설토토 입금내역 유니벳 Games, Casino, Games, 스포츠 토토 축구 승무패 벳피스트 놀검소 Casino, Games, Casino, Games, Casino,
bergama transfer
回覆刪除güzelbahçe transfer
foça transfer
mordoğan transfer
aliağa transfer
balıkesir transfer
7YOWQ