Linux 环境下 intel 核心显卡的 OpenCL 并行计算 - Chinaunix

// Create sobel filter program
extern const std::string str_cl_sobel = BOOST_COMPUTE_STRINGIZE_SOURCE
(
//For out of boundary pixels, edge pixel
// value is returned
const sampler_t sampler = CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_NEAREST;
kernel void cl_kernel(read_only image2d_t src, write_only image2d_t dst)
{
int x = (int)get_global_id(0);
int y = (int)get_global_id(1);
if (x >= get_image_width(src) || y >= get_image_height(src))
return;
// [(x-1, y+1), (x, y+1), (x+1, y+1)]
// [(x-1, y ), (x, y ), (x+1, y )]
// [(x-1, y-1), (x, y-1), (x+1, y-1)]
// [p02, p12, p22]
// [p01, pixel, p21]
// [p00, p10, p20]
//Basically finding influence of neighbour pixels on current pixel
float4 p00 = read_imagef(src, sampler, (int2)(x - 1, y - 1));
float4 p10 = read_imagef(src, sampler, (int2)(x, y - 1));
float4 p20 = read_imagef(src, sampler, (int2)(x + 1, y - 1));
float4 p01 = read_imagef(src, sampler, (int2)(x - 1, y));
//pixel that we are working on
float4 p21 = read_imagef(src, sampler, (int2)(x + 1, y));
float4 p02 = read_imagef(src, sampler, (int2)(x - 1, y + 1));
float4 p12 = read_imagef(src, sampler, (int2)(x, y + 1));
float4 p22 = read_imagef(src, sampler, (int2)(x + 1, y + 1));
//Find Gx = kernel + 3x3 around current pixel
// Gx = [-1 0 +1] [p02, p12, p22]
// [-2 0 +2] + [p01, pixel, p21]
// [-1 0 +1] [p00, p10, p20]
float3 gx = -p00.xyz + p20.xyz +
2.0f * (p21.xyz - p01.xyz)
-p02.xyz + p22.xyz;
//Find Gy = kernel + 3x3 around current pixel
// Gy = [-1 -2 -1] [p02, p12, p22]
// [ 0 0 0] + [p01, pixel, p21]
// [+1 +2 +1] [p00, p10, p20]
float3 gy = p00.xyz + p20.xyz +
2.0f * (- p12.xyz + p10.xyz) -
p02.xyz - p22.xyz;
//Find G
float3 g = native_sqrt(gx * gx + gy * gy);
// we could also approximate this as g = fabs(gx) + fabs(gy)
write_imagef(dst, (int2)(x, y), (float4)(g.x, g.y, g.z, 1.0f));
}
);

复制代码