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 a7622f
//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 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 a7622f
	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
	float2 d;
Ivan Mahonin 6fa009
	d.x = p1.x - p0.x;
Ivan Mahonin 6fa009
	d.y = p1.y - p0.y;
Ivan Mahonin 6fa009
	int w1 = width - 1;
Ivan Mahonin 6fa009
	int h1 = height - 1;
Ivan Mahonin a7622f
	float kx = d.x/d.y;
Ivan Mahonin a7622f
	float ky = d.y/d.x;
Ivan Mahonin 6fa009
	
Ivan Mahonin 6fa009
	while(p0.x != p1.x || p0.y != p1.y) {
Ivan Mahonin a7622f
		int ix = max((int)p0.x, 0);
Ivan Mahonin a7622f
		int iy = (int)p0.y;
Ivan Mahonin a7622f
		if (ix > w1) 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 a7622f
		iy = max(0, min(h1, iy));
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 a7622f
		
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 a7622f
		float cover = (pp1.x - p0.x)*ONE_F;
Ivan Mahonin a7622f
		float area = py.y - 0.5f*(p0.y + pp1.y);
Ivan Mahonin a7622f
		if (flipx) { ix = w1 - ix; cover = -cover; }
Ivan Mahonin a7622f
		if (flipy) { iy = h1 - iy; area = 1.f - area; }
Ivan Mahonin a7622f
		p0 = pp1;
Ivan Mahonin 6fa009
		
Ivan Mahonin 6fa009
		atomicAdd(
Ivan Mahonin a7622f
			(unsigned long long*)(marks + 2*(iy*width + 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 a7622f
		//int *mark = marks + ((iy*width + ix) << 1);
Ivan Mahonin a7622f
		//atomicAdd(mark, (int)(area*cover));
Ivan Mahonin a7622f
		//atomicAdd(mark + 1, (int)(cover));
Ivan Mahonin 6fa009
	}
Ivan Mahonin 6fa009
}
Ivan Mahonin 6fa009
Ivan Mahonin 6fa009
__global__ void fill(
Ivan Mahonin 6fa009
	int width,
Ivan Mahonin a7622f
	int2 *marks,
Ivan Mahonin 6fa009
	float4 *image,
Ivan Mahonin 6fa009
	float4 color,
Ivan Mahonin 6fa009
	int4 bounds )
Ivan Mahonin 6fa009
{
Ivan Mahonin a7622f
	int id = blockIdx.x*blockDim.x + threadIdx.x + bounds.x;
Ivan Mahonin a7622f
	if (id >= bounds.z) return;
Ivan Mahonin a7622f
	id += bounds.y*width;
Ivan Mahonin 6fa009
	marks += id;
Ivan Mahonin 6fa009
	image += id;
Ivan Mahonin 6fa009
Ivan Mahonin a7622f
	int icover = 0;
Ivan Mahonin a7622f
	while(true) {
Ivan Mahonin a7622f
		int2 m = *marks;
Ivan Mahonin a7622f
		*marks = make_int2(0, 0);
Ivan Mahonin a7622f
		float alpha = (float)abs(m.x + icover)*color.w*DIV_ONE_F;
Ivan Mahonin a7622f
		marks += width;
Ivan Mahonin 6fa009
Ivan Mahonin a7622f
		icover += m.y;
Ivan Mahonin 6fa009
		float one_alpha = 1.f - alpha;
Ivan Mahonin 6fa009
		
Ivan Mahonin a7622f
		float4 p = *image;
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 a7622f
		*image = p;
Ivan Mahonin 6fa009
		
Ivan Mahonin a7622f
		if (++bounds.y >= bounds.w) return;
Ivan Mahonin a7622f
		image += width;
Ivan Mahonin 6fa009
	}
Ivan Mahonin 6fa009
}
Ivan Mahonin 6fa009
Ivan Mahonin a7622f
//}