Blame c++/contourgl/cl/contour.cl

d989ab
/*
d989ab
    ......... 2015 Ivan Mahonin
d989ab
d989ab
    This program is free software: you can redistribute it and/or modify
d989ab
    it under the terms of the GNU General Public License as published by
d989ab
    the Free Software Foundation, either version 3 of the License, or
d989ab
    (at your option) any later version.
d989ab
d989ab
    This program is distributed in the hope that it will be useful,
d989ab
    but WITHOUT ANY WARRANTY; without even the implied warranty of
d989ab
    MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
d989ab
    GNU General Public License for more details.
d989ab
d989ab
    You should have received a copy of the GNU General Public License
d989ab
    along with this program.  If not, see <http: licenses="" www.gnu.org="">.</http:>
d989ab
*/
d989ab
013f0c
// paths format:
013f0c
//   {
013f0c
//     int count,
013f0c
//     paths: [
013f0c
//       {
013f0c
//         int point_count,
013f0c
//         int flags,
013f0c
//         float4 color,
013f0c
//         points: [ float2, ... ]
013f0c
//       },
013f0c
//       ...
013f0c
//     ]
013f0c
//   }
013f0c
67876c
013f0c
kernel void draw(
013f0c
	global char *paths_buffer,
c7fa36
	global int *mark_buffer,
013f0c
	read_only image2d_t read_image,
013f0c
	write_only image2d_t write_image ) // assumed that read and write image is the same object
d989ab
{
5890eb
	const float e = 1e-6f;
013f0c
013f0c
	int id = (int)get_global_id(0);
013f0c
	int count = (int)get_global_size(0);
013f0c
	
013f0c
	int paths_count = *(global int *)paths_buffer;
013f0c
	global char *paths = paths_buffer + sizeof(int);
c7fa36
	
013f0c
	int width = get_image_width(write_image);
013f0c
	int height = get_image_height(write_image);
013f0c
	int pixels_count = width*height;
013f0c
	float2 size = (float2)((float)width, (float)height); 
c7fa36
	int w1 = width - 1;
c7fa36
	int h1 = height - 1;
c7fa36
	
013f0c
	global int *bound_minx = (global int *)(mark_buffer + 2*pixels_count);
013f0c
	global int *bound_miny = bound_minx + 1;
013f0c
	global int *bound_maxx = bound_minx + 2;
013f0c
	global int *bound_maxy = bound_minx + 3;
013f0c
	
013f0c
	// clear marks
013f0c
	for(int i = id; i < 2*pixels_count; i += count)
013f0c
		mark_buffer[i] = 0;
013f0c
	barrier(CLK_LOCAL_MEM_FENCE);
013f0c
	
013f0c
	// draw paths
013f0c
	for(int p = 0; p < paths_count; ++p) {
013f0c
		int points_count = *(global int *)paths; paths += sizeof(int);
013f0c
		int flags        = *(global int *)paths; paths += sizeof(int);
c7fa36
		
013f0c
		float4 color;
013f0c
		color.x = *(global float *)paths; paths += sizeof(float);
013f0c
		color.y = *(global float *)paths; paths += sizeof(float);
013f0c
		color.z = *(global float *)paths; paths += sizeof(float);
013f0c
		color.w = *(global float *)paths; paths += sizeof(float);
c7fa36
		
013f0c
		global float *points = (global float *)paths;
013f0c
		paths += 2*points_count*sizeof(float);
013f0c
		
013f0c
		int segments_count = points_count - 1;
013f0c
		if (segments_count <= 0) continue;
013f0c
		
013f0c
		int invert  = flags & 1;
013f0c
		int evenodd = flags & 2;
013f0c
		
013f0c
		*bound_minx = invert ?  0 : (int)floor(clamp(points[0] + e, 0.f, size.x - 1.f + e));
013f0c
		*bound_miny = invert ?  0 : (int)floor(clamp(points[1] + e, 0.f, size.y - 1.f + e));
013f0c
		*bound_maxx = invert ? w1 : *bound_minx;
013f0c
		*bound_maxy = invert ? h1 : *bound_miny;
d989ab
013f0c
		// trace path
013f0c
		for(int i = id; i < segments_count; i += count) {
013f0c
			int ii = 2*i;
013f0c
			float2 p0 = { points[ii + 0], points[ii + 1] };
013f0c
			float2 p1 = { points[ii + 2], points[ii + 3] };
013f0c
			
013f0c
			int p1x = (int)floor(clamp(p1.x + e, 0.f, size.x - 1.f + e));
013f0c
			int p1y = (int)floor(clamp(p1.y + e, 0.f, size.y - 1.f + e));
013f0c
			atomic_min(bound_minx, p1x - 1);
013f0c
			atomic_min(bound_miny, p1y - 1);
013f0c
			atomic_max(bound_maxx, p1x + 1);
013f0c
			atomic_max(bound_maxy, p1y + 1);
013f0c
				
013f0c
			bool flipx = p1.x < p0.x;
013f0c
			bool flipy = p1.y < p0.y;
013f0c
			if (flipx) { p0.x = size.x - p0.x; p1.x = size.x - p1.x; }
013f0c
			if (flipy) { p0.y = size.y - p0.y; p1.y = size.y - p1.y; }
013f0c
			float2 d = p1 - p0;
013f0c
			float kx = fabs(d.y) < e ? 1e10 : d.x/d.y;
013f0c
			float ky = fabs(d.x) < e ? 1e10 : d.y/d.x;
013f0c
			
013f0c
			while(p0.x != p1.x || p0.y != p1.y) {
013f0c
				int ix = (int)floor(p0.x + e);
013f0c
				int iy = (int)floor(p0.y + e);
013f0c
				if (iy > h1) break;
0b3288
013f0c
				float2 px, py;
013f0c
				px.x = (float)(ix + 1);
013f0c
				px.y = p0.y + ky*(px.x - p0.x);
013f0c
				py.y = max((float)(iy + 1), 0.f);
013f0c
				py.x = p0.x + kx*(py.y - p0.y);
013f0c
				float2 pp1 = p1;
013f0c
				if (pp1.x > px.x) pp1 = px;
013f0c
				if (pp1.y > py.y) pp1 = py;
013f0c
				
013f0c
				if (iy >= 0) {
013f0c
					// calc values
013f0c
					float cover = pp1.y - p0.y;
013f0c
					float area = px.x - 0.5f*(p0.x + pp1.x);
013f0c
					if (flipx) { ix = w1 - ix; area = 1.f - area; }
013f0c
					if (flipy) { iy = h1 - iy; cover = -cover; }
013f0c
					ix = clamp(ix, 0, w1);
013f0c
					
013f0c
					// store in buffer
013f0c
					global int *mark = mark_buffer + (iy*width + ix)*2;
013f0c
					atomic_add(mark, (int)round(area*cover*65536.f));
013f0c
					atomic_add(mark + 1, (int)round(cover*65536.f));
013f0c
				}
013f0c
				
013f0c
				p0 = pp1;
013f0c
			}
013f0c
		}
013f0c
		barrier(CLK_LOCAL_MEM_FENCE);
a04770
		
013f0c
		// fill
013f0c
		int2 coord;
013f0c
		int minx = max(*bound_minx, 0);
013f0c
		int miny = max(*bound_miny, 0);
013f0c
		int maxx = min(*bound_maxx, w1);
013f0c
		int maxy = min(*bound_maxy, h1);
013f0c
		for(coord.y = miny + id; coord.y <= maxy; coord.y += count) {
013f0c
			global int *mark = mark_buffer + (coord.y*width + minx)*2;
013f0c
	
013f0c
			float cover = 0.f;
013f0c
			for(coord.x = minx; coord.x <= maxx; ++coord.x) {
013f0c
				// read mark (alpha, cover)
013f0c
				float alpha = fabs(cover + *mark/65536.f); *mark = 0; ++mark;
013f0c
				cover += *mark/65536.f;                    *mark = 0; ++mark;
013f0c
				
013f0c
				//if (evenodd) alpha = 1.f - fabs(fmod(alpha, 2.f) - 1.f);
013f0c
				//if (invert) alpha = 1.f - alpha;
013f0c
				alpha *= color.w;
013f0c
				
013f0c
				// write color
013f0c
				float alpha_inv = 1.f - alpha;
013f0c
				float4 cl = read_imagef(read_image, coord);
013f0c
				cl.x = cl.x*alpha_inv + color.x*alpha;
013f0c
				cl.y = cl.y*alpha_inv + color.y*alpha;
013f0c
				cl.z = cl.z*alpha_inv + color.z*alpha;
013f0c
				cl.w = min(cl.w + alpha, 1.f);
013f0c
				write_imagef(write_image, coord, cl);
013f0c
			}
013f0c
		}
013f0c
		barrier(CLK_LOCAL_MEM_FENCE);
d989ab
	}
d989ab
}
013f0c