double2 sqrt( const expression<E1>& a, double2_type )
{
double2 w,x;
w = a();
x = native_sqrt(w,double2_type());
x = 0.5*(x + w/x);
x = 0.5*(x + w/x);
return x;
}
kernel void sobel_rgb(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));
}
double2 native_sqrt( const expression<E1>& e1, double2_type )
{
typedef unary<E1,cal_unary_sqrt<typename E1::value_type> > sqrt_type;
#if defined(__CAL_H__)
if( Source::info().available && Source::info().target>=CAL_TARGET_CYPRESS ) return sqrt_type(e1());
#endif
double2 w,x;
int2 e;
w = e1();
x = frexp(w,e);
x = cast_type<double2>( native_sqrt(cast_type<float2>(x),float2_type()) );
x = select( e&0x1, 1.41421356237309504880*x, x );
x = ldexp(x,e>>1);
return x;
}
double1 fast_sqrt( const expression<E1>& a, float4_type )
{
return native_sqrt(a(),float4_type());
}
