Blame c++/contourgl/cuda/contour.cu

6fa009
/*
6fa009
    ......... 2018 Ivan Mahonin
6fa009
6fa009
    This program is free software: you can redistribute it and/or modify
6fa009
    it under the terms of the GNU General Public License as published by
6fa009
    the Free Software Foundation, either version 3 of the License, or
6fa009
    (at your option) any later version.
6fa009
6fa009
    This program is distributed in the hope that it will be useful,
6fa009
    but WITHOUT ANY WARRANTY; without even the implied warranty of
6fa009
    MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
6fa009
    GNU General Public License for more details.
6fa009
6fa009
    You should have received a copy of the GNU General Public License
6fa009
    along with this program.  If not, see <http: licenses="" www.gnu.org="">.</http:>
6fa009
*/
6fa009
6fa009
extern "C" {
6fa009
6fa009
#define ONE       65536
6fa009
#define TWO      131072               // (ONE)*2
6fa009
#define HALF      32768               // (ONE)/2
6fa009
#define ONE_F     65536.f             // (float)(ONE)
6fa009
#define DIV_ONE_F 0.0000152587890625f // 1.f/(ONE_F)
6fa009
6fa009
6fa009
__global__ void clear(
6fa009
	int width,
6fa009
	int height,
6fa009
	int4 *marks )
6fa009
{
6fa009
	int id = blockIdx.x*blockDim.x + threadIdx.x;
6fa009
	int c = id % width;
6fa009
	marks[id] = make_int4(0, 0, c | (c + 1), 0);
6fa009
}
6fa009
6fa009
__global__ void path(
6fa009
	int width,
6fa009
	int height,
6fa009
	int *marks,
6fa009
	const float2 *points,
6fa009
	int begin,
6fa009
	int end,
6fa009
	int minx )
6fa009
{
6fa009
	int id = blockIdx.x*blockDim.x + threadIdx.x + begin;
6fa009
	if (id >= end) return;
6fa009
	float2 p0 = points[id];
6fa009
	float2 p1 = points[id + 1];
6fa009
	
6fa009
	bool flipx = p1.x < p0.x;
6fa009
	bool flipy = p1.y < p0.y;
6fa009
	if (flipx) { p0.x = (float)width - p0.x; p1.x = (float)width - p1.x; }
6fa009
	if (flipy) { p0.y = (float)height - p0.y; p1.y = (float)height - p1.y; }
6fa009
	
6fa009
	float2 d;
6fa009
	d.x = p1.x - p0.x;
6fa009
	d.y = p1.y - p0.y;
6fa009
	float kx = d.x/d.y;
6fa009
	float ky = d.y/d.x;
6fa009
	int w1 = width - 1;
6fa009
	int h1 = height - 1;
6fa009
	
6fa009
	while(p0.x != p1.x || p0.y != p1.y) {
6fa009
		int ix = (int)p0.x;
6fa009
		int iy = max((int)p0.y, 0);
6fa009
		if (iy > h1) return;
6fa009
6fa009
		float2 px, py;
6fa009
		px.x = (float)(ix + 1);
6fa009
		py.y = (float)(iy + 1);
6fa009
		ix = max(0, min(w1, ix));
6fa009
		
6fa009
		px.y = p0.y + ky*(px.x - p0.x);
6fa009
		py.x = p0.x + kx*(py.y - p0.y);
6fa009
6fa009
		float2 pp1 = p1;
6fa009
		if (pp1.x > px.x) pp1 = px;
6fa009
		if (pp1.y > py.y) pp1 = py;
6fa009
		
6fa009
		float cover = (pp1.y - p0.y)*ONE_F;
6fa009
		float area = px.x - 0.5f*(p0.x + pp1.x);
6fa009
		if (flipx) { ix = w1 - ix; area = 1.f - area; }
6fa009
		if (flipy) { iy = h1 - iy; cover = -cover; }
6fa009
		
6fa009
		int *row = marks + 4*iy*width;
6fa009
		atomicAdd(
6fa009
			(unsigned long long*)(row + 4*ix),
6fa009
			((unsigned long long)(unsigned int)(int)(cover) << 32)
6fa009
			| (unsigned long long)(unsigned int)((int)(area*cover)) );
6fa009
		//row[4*ix] += (int)(area*cover);
6fa009
		//row[4*ix + 1] += (int)(cover);
6fa009
		//atomicAdd(row + 4*ix, (int)(area*cover));
6fa009
		//atomicAdd(row + 4*ix + 1, (int)(cover));
6fa009
		
6fa009
		row += 2;
6fa009
		int iix = (ix & (ix + 1)) - 1;
6fa009
		while(iix >= minx) {
6fa009
			atomicMin(row + 4*iix, ix);
6fa009
			iix = (iix & (iix + 1)) - 1;
6fa009
		}
6fa009
		
6fa009
		p0 = pp1;
6fa009
	}
6fa009
}
6fa009
6fa009
__global__ void fill(
6fa009
	int width,
6fa009
	int4 *marks,
6fa009
	float4 *image,
6fa009
	float4 color,
6fa009
	int4 bounds )
6fa009
{
6fa009
	int id = blockIdx.x*blockDim.x + threadIdx.x + bounds.y;
6fa009
	if (id >= bounds.w) return;
6fa009
	id *= width;
6fa009
	marks += id;
6fa009
	image += id;
6fa009
	
6fa009
	int4 *mark;
6fa009
	float4 *pixel;
6fa009
6fa009
	int4 m;
6fa009
	int icover = 0, c0 = bounds.x, c1 = bounds.x;
6fa009
	while(c1 < bounds.z) {
6fa009
		if (abs(icover) > HALF)
6fa009
			while(c0 < c1)
6fa009
				image[c0++] = color;
6fa009
6fa009
		mark = &marks[c1];
6fa009
		m = *mark;
6fa009
		*mark = make_int4(0, 0, c1 | (c1 + 1), 0); 
6fa009
		
6fa009
		float alpha = (float)abs(m.x + icover)*DIV_ONE_F;
6fa009
		float one_alpha = 1.f - alpha;
6fa009
		
6fa009
		pixel = &image[c1];
6fa009
		float4 p = *pixel;
6fa009
		p.x = p.x*one_alpha + color.x*alpha;
6fa009
		p.y = p.y*one_alpha + color.y*alpha;
6fa009
		p.z = p.z*one_alpha + color.z*alpha;
6fa009
		p.w = p.w*one_alpha + color.w*alpha;
6fa009
		*pixel = p;
6fa009
		
6fa009
		icover += m.y;
6fa009
		c0 = c1 + 1;
6fa009
		c1 = m.z;
6fa009
	}
6fa009
}
6fa009
6fa009
}