openclbasetransform: Fix some bug that raises a segfault.
[valastuff:gst-plugins-cl.git] / src / example_kernels.cl
1
2 /*
3  * A sobel filter, original taken from http://code.imagej.net/trac/imagej/browser/trunk/opencl/sobel.cl
4  */
5 __kernel void sobel( __global uchar* output, __global uchar* input, int width, int height)
6 {
7     int x = get_global_id(0);
8     int y = get_global_id(1);
9     int p[9];
10     int offset = y * width + x;
11 if( x < 1 || y < 1 || x > width - 2 || y > height - 2 )
12 {
13   output[offset] = 0; //TODO implement image edges
14 }
15 else
16 {
17     p[0] = input[offset - width - 1] & 0xff;
18     p[1] = input[offset - width] & 0xff;
19     p[2] = input[offset - width + 1] & 0xff;
20     p[3] = input[offset - 1] & 0xff;
21     p[4] = input[offset] & 0xff;
22     p[5] = input[offset + 1] & 0xff;
23     p[6] = input[offset + width - 1] & 0xff;
24     p[7] = input[offset + width] & 0xff;
25     p[8] = input[offset + width + 1] & 0xff;
26
27     int sum1 = p[0] + 2*p[1] + p[2] - p[6] - 2*p[7] - p[8];
28     int sum2 = p[0] + 2*p[3] + p[6] - p[2] - 2*p[5] - p[8];
29     float sum3 = sum1*sum1 + sum2*sum2;
30    
31     int sum = sqrt( sum3 );
32     if (sum > 255) sum = 255;
33     output[offset] = (char) sum;
34  }
35 };
36
37
38 #define BI_RADIX 20
39 __kernel void 
40 bidiff (__global       uchar* dst, 
41         __global const uchar* src, 
42                  const int width,
43                  const int height)
44 {
45   int x = get_global_id (0),
46       y = get_global_id (1);
47   
48   int idx = y * width + x;
49   int w = width / 2;
50   
51   if (x < w)
52   {
53     int n = 0;
54     float r = 0;
55     float sum_a = 0, 
56           sum_b = 0;
57     
58     for (int lx = clamp(x-BI_RADIX, 0, w) ; lx < clamp(x+BI_RADIX, 0, w) ; lx++)
59     {
60       for (int ly = clamp(y-BI_RADIX, 0, height) ; ly < clamp(y+BI_RADIX, 0, height) ; ly++)
61       {
62         int lidx = ly * width + lx;
63         float lv = convert_float(src[lidx]), 
64               rv = convert_float(src[lidx + w]);
65
66         r += lv * rv;
67         sum_a += lv * lv;
68         sum_b += rv * rv;
69         n++;
70       }
71     }
72     r = r / (sqrt(sum_a) * sqrt(sum_b));
73     //r = r / (n*255*255);
74     if (r > 1 || r < 0)
75     {
76       dst[idx+w] = 0xff;
77     }
78     else
79     {
80       r *= 255;
81     }
82     dst[idx] = (r);
83   }
84   else
85   {
86     dst[idx] = src[idx];
87     dst[idx] = abs(src[idx] - src[idx+w]);
88   }
89 }
90
91 __kernel void 
92 imrotate (__write_only  image2d_t dst, 
93           __read_only   image2d_t src, 
94           const         sampler_t src_sampler, 
95           const         int       width,
96           const         int       height)
97 {
98   const int x = get_global_id (0);
99   const int y = get_global_id (1);
100   /*int x2 = 0, y2 = 0;
101   float a = 0.1;
102   
103   float fx = x, 
104         fy = y;
105   
106   x2 = fx * cos(a) + fy * sin(a);
107   y2 = -fx * sin(a) + fy * cos(a);
108   */
109   uint4 val = read_imageui (src, src_sampler, (int2) (x, y));
110   write_imageui(dst, (int2)( x, y ), val);
111 }