__kernel void kern(
__read_only image2d_t entry,
__read_only image2d_t exit,
__write_only image2d_t tex,
__read_only image1d_t transfer,
__global float * data,
int width,
int height,
int depth
)
{
int x = get_global_id(0);
int y = get_global_id(1);
int2 coords = (int2)(x,y);
float2 tcoords = (float2)(x,y)/512.0f;
const float Samplings = 100.0f;
const float k = 3.0f;
float3 a=read_imagef(entry,samplersrc,tcoords).xyz;
float3 b=read_imagef(exit,samplersrc,tcoords).xyz;
float3 dir=b-a;
int steps = (int)(floor(Samplings * length(dir)));
float3 diff1 = dir / (float)(steps);
dir=dir*(1.0f/length(dir));
float delta=1.0f/Samplings;
float4 result = (float4)(0.0f,0.0f,0.0f,1.0f);
for (int i=0; i<steps; i++) {
float3 p=a;
float valuex = 0.0f;
float3 valued = 0.0f;
//Calculate gradients
evaldx(data, p.x, p.y, p.z, width, height, depth, &valuex, &valued);
//Apply classification
float4 color=read_imagef(transfer,samplersrc,valuex);
result.w*=pow(color.w,delta);
result.xyz+=result.w*color.xyz*delta;
if(result.w<0.05f) {
i=steps;
result.w=0.0f;
}
a+=diff1;
}
write_imagef(tex, coords, result);
};
__kernel void merge(__write_only image2d_t destination, __read_only image2d_t previousDestination, __read_only image2d_t sourceA, __read_only image2d_t sourceB, int xA, int yA, int xB, int yB)
{
int2 destinationCoord = (int2) (get_global_id(0), get_global_id(1));
int2 sourceCoordA = (int2) (destinationCoord.x + xA, destinationCoord.y + yA);
int2 sourceCoordB = (int2) (destinationCoord.x + xB, destinationCoord.y + yB);
float4 destinationPixel = read_imagef(previousDestination, sampler, destinationCoord);
float4 sourcePixelA = read_imagef(sourceA, sampler, sourceCoordA);
float4 sourcePixelB = read_imagef(sourceB, sampler, sourceCoordB);
destinationPixel = sourcePixelA + destinationPixel * (1 - sourcePixelA.w);
destinationPixel = sourcePixelB + destinationPixel * (1 - sourcePixelB.w);
write_imagef(destination, destinationCoord, destinationPixel);
}
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));
}
__kernel void resize_image(__read_only image2d_t input,
const sampler_t sampler,
__write_only image2d_t output)
{
const uint x = get_global_id(0);
const uint y = get_global_id(1);
const float w = get_image_width(output);
const float h = get_image_height(output);
float2 coord = { ((float) x / w) * get_image_width(input),
((float) y / h) * get_image_height(input) };
float4 pixel = read_imagef(input, sampler, coord);
write_imagef(output, (int2)(x, h - y - 1), pixel);
};
__kernel void optical_flow (
read_only
image2d_t current_image,
read_only image2d_t previous_image,
write_only image2d_t optical_flow,
const float scale,
const float offset,
const float lambda,
const float threshold )
{
sampler_t sampler = CLK_ADDRESS_CLAMP_TO_EDGE;
int2 coords = (int2)(get_global_id(0), get_global_id(1));
float4 current_pixel = read_imagef(current_image,
sampler,
coords);
float4 previous_pixel = read_imagef(previous_image,
sampler,
coords);
int2 x1 = (int2)(offset, 0.f);
int2 y1 = (int2)(0.f, offset);
//get the difference
float4 curdif = previous_pixel - current_pixel;
//calculate the gradient
//Image 2 first
float4 gradx = read_imagef(previous_image,
sampler,
coords+x1) -
read_imagef(previous_image,
sampler,
coords-x1);
//Image 1
gradx += read_imagef(current_image,
sampler,
coords+x1) -
read_imagef(current_image,
sampler,
coords-x1);
//Image 2 first
float4 grady = read_imagef(previous_image,
sampler,
coords+y1) -
read_imagef(previous_image,
sampler,
coords-y1);
//Image 1
grady += read_imagef(current_image,
sampler,
coords+y1) -
read_imagef(current_image,
sampler,
coords-y1);
float4 sqr = (gradx*gradx) + (grady*grady) +
(float4)(lambda,lambda, lambda, lambda);
float4 gradmag = sqrt(sqr);
///////////////////////////////////////////////////
float4 vx = curdif * (gradx / gradmag);
float vxd = vx.x;//assumes greyscale
//format output for flowrepos, out(-x,+x,-y,+y)
float2 xout = (float2)(fmax(vxd,0.f),fabs(fmin(vxd,0.f)));
xout *= scale;
///////////////////////////////////////////////////
float4 vy = curdif*(grady/gradmag);
float vyd = vy.x;//assumes greyscale
//format output for flowrepos, out(-x,+x,-y,+y)
float2 yout = (float2)(fmax(vyd,0.f),fabs(fmin(vyd,0.f)));
yout *= scale;
///////////////////////////////////////////////////
float4 out = (float4)(xout, yout);
float cond = (float)isgreaterequal(length(out), threshold);
out *= cond;
write_imagef(optical_flow, coords, out);
}
__kernel void convolution(__read_only image2d_t sourceImage,
__write_only image2d_t outputImage,
__constant float* filter,
int filterWidth)
{
const sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE |
CLK_ADDRESS_CLAMP_TO_EDGE |
CLK_FILTER_NEAREST;
// Store each work-item’s unique row and column
int x = get_global_id(0);
int y = get_global_id(1);
// Half the width of the filter is needed for indexing
// memory later
int halfWidth = (int)(filterWidth/2);
// All accesses to images return data as four-element vector
// (i.e., float4).
float4 sum = {0.0f, 0.0f, 0.0f, 0.0f};
// Iterator for the filter
int filterIdx = 0;
// Each work-item iterates around its local area based on the
// size of the filter
int2 coords; // Coordinates for accessing the image
// Iterate the filter rows
for(int i = -halfWidth; i <= halfWidth; i++)
{
coords.y = y + i;
// Iterate over the filter columns
for(int j = -halfWidth; j <= halfWidth; j++)
{
coords.x = x + j;
float4 pixel;
// Read a pixel from the image.
// Work on a channel
pixel = read_imagef(sourceImage, sampler, coords);
sum.x += pixel.x * filter[filterIdx++];
//sum.y += pixel.y * filter[filterIdx++];
//sum.z += pixel.z * filter[filterIdx++];
}
}
barrier(CLK_GLOBAL_MEM_FENCE);
// Copy the data to the output image if the
// work-item is in bounds
if(y < get_image_height(sourceImage) &&
x < get_image_width(sourceImage))
{
coords.x = x;
coords.y = y;
//Same channel is copied in all three channels
//write_imagef(outputImage, coords,
// (float4)(sum.x,sum.x,sum.x,1.0f));
write_imagef(outputImage, coords, sum);
}
}
