|
|
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 |
}
|