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

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