diff --git a/c++/contourgl/Makefile b/c++/contourgl/Makefile index 98eca6f..ba6608e 100644 --- a/c++/contourgl/Makefile +++ b/c++/contourgl/Makefile @@ -1,6 +1,6 @@ DEPLIBS = gl x11 OpenCL -CXXFLAGS = -O3 -Wall -fmessage-length=0 `pkg-config --cflags $(DEPLIBS)` -DGL_GLEXT_PROTOTYPES +CXXFLAGS = -O0 -g -Wall -fmessage-length=0 `pkg-config --cflags $(DEPLIBS)` -DGL_GLEXT_PROTOTYPES HEADERS = \ clcontext.h \ diff --git a/c++/contourgl/cl/contour.cl b/c++/contourgl/cl/contour.cl new file mode 100644 index 0000000..bbe9324 --- /dev/null +++ b/c++/contourgl/cl/contour.cl @@ -0,0 +1,131 @@ +/* + ......... 2015 Ivan Mahonin + + This program is free software: you can redistribute it and/or modify + it under the terms of the GNU General Public License as published by + the Free Software Foundation, either version 3 of the License, or + (at your option) any later version. + + This program is distributed in the hope that it will be useful, + but WITHOUT ANY WARRANTY; without even the implied warranty of + MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the + GNU General Public License for more details. + + You should have received a copy of the GNU General Public License + along with this program. If not, see . +*/ + +__kernel void lines( + __global int width, + __global float *lines, + __global int *rows, + __global float *mark_buffer ) +{ + size_t id = get_global_id(0); + int *row = rows + id*2; + int begin = *rows; + int end = begin + rows[1]; + for(int i = begin; i < end; ++i) { + float *line = lines + 4*begin; + float2 p0(*line, line[1]); + float2 p1(line[2], line[3]); + + int iy0 = (int)floor(p0.y); + int iy1 = (int)floor(p1.x); + if (iy1 < iy0) { int sw = iy0; iy0 = iy1; iy1 = iy0; } + + float2 d = p1 - p0; + float2 k( fabs(d.y) < 1e-6 ? 0.0 : d.x/d.y, + fabs(d.x) < 1e-6 ? 0.0 : d.y/d.x ); + + for(int r = iy0; r <= iy1; ++r) { + float y = (float)iy0; + + float2 pp0 = p0; + pp0.y -= y; + if (pp0.y < 0.0) { + pp0.y = 0.0; + pp0.x = p0.x - k.x*y; + } else + if (pp0.y > 1.0) { + pp0.y = 1.0; + pp0.x = p0.x - k.x*(y - 1.0); + } + + float2 pp1 = p1; + pp1.y -= y; + if (pp1.y < 0.0) { + pp1.y = 0.0; + pp1.x = p0.x - k.x*y; + } else + if (pp1.y > 1.0) { + pp1.y = 1.0; + pp1.x = p0.x - k.x*(y - 1.0); + } + + int ix0 = min(max((int)floor(pp0.x), 0), width); + int ix1 = min(max((int)floor(pp1.x), 0), width); + if (ix1 < ix0) { int sw = ix0; ix0 = ix1; ix1 = ix0; } + for(int c = ix0; c <= ix1; ++c) { + float x = (float)ix0; + + float2 ppp0 = pp0; + ppp0.x -= x; + if (ppp0.x < 0.0) { + ppp0.x = 0.0; + ppp0.y = pp0.y - k.y*x; + } else + if (ppp0.x > 1.0) { + ppp0.x = 1.0; + ppp0.y = pp0.y - k.y*(x - 1.0); + } + + float2 ppp1 = pp1; + ppp1.x -= x; + if (ppp1.x < 0.0) { + ppp1.x = 0.0; + ppp1.y = pp0.y - k.y*x; + } else + if (ppp1.x > 1.0) { + ppp1.x = 1.0; + ppp1.y = pp0.y - k.y*(x - 1.0); + } + + float cover = ppp0.y - ppp1.y; + float area = (0.5*(ppp1.x + ppp1.x) - 1.0)*cover; + float *mark = mark_buffer + 2*(r*width + c); + *mark += area; + mark[1] += cover; + } + } + } +} + +__kernel void fill( + __global int width, + __global float *mark_buffer, + __global float *surface_buffer, + __global float color_r, + __global float color_g, + __global float color_b, + __global float color_a ) +{ + sizet id = get_global_id(0); + int w = width; + float cr = color_r; + float cg = color_g; + float cb = color_b; + float ca = color_a; + float *mark = mark_buffer + 2*id*width; + float *surface = surface_buffer + 4*id*width; + float cover = 0; + for(int i = 0; i < w; ++i, mark += 2, surface += 4) { + float alpha = ca*(1.0 - fabs(1.0 - 0.5*frac(fabs(2.0*(*mark + cover))))); + float alpha_inv = 1.0 - alpha; + surface[0] = surface[0]*alpha_inv + cr*alpha; + surface[1] = surface[1]*alpha_inv + cg*alpha; + surface[2] = surface[2]*alpha_inv + cb*alpha; + surface[3] = surface[3]*alpha_inv + ca*alpha; + cover += mark[1]; + } +} diff --git a/c++/contourgl/cl/hello.cl b/c++/contourgl/cl/hello.cl index b93cfc1..a361161 100644 --- a/c++/contourgl/cl/hello.cl +++ b/c++/contourgl/cl/hello.cl @@ -1,6 +1,6 @@ __constant char s[] = "Hello!"; __kernel void hello(__global char *out) { - size_t tid = get_global_id(0); - out[tid] = s[tid]; + size_t i = get_global_id(0); + out[i] = s[i]; } \ No newline at end of file diff --git a/c++/contourgl/clrender.cpp b/c++/contourgl/clrender.cpp index d769b08..b0d8b0a 100644 --- a/c++/contourgl/clrender.cpp +++ b/c++/contourgl/clrender.cpp @@ -120,7 +120,7 @@ Surface* ClRender::receive_surface() { } -void ClRender::contour(const Contour &contour, const Rect &rect, const Color &color) { +void ClRender::contour(const Contour &contour, const Rect &rect, const Color &color, bool invert, bool evenodd) { Measure t("ClRender::contour"); Contour transformed, splitted; @@ -138,6 +138,7 @@ void ClRender::contour(const Contour &contour, const Rect &rect, const Color &co { Measure t("split"); + transformed.allow_split_lines = true; transformed.split(splitted, to, Vector(0.5, 0.5)); } @@ -206,18 +207,22 @@ void ClRender::contour(const Contour &contour, const Rect &rect, const Color &co clFinish(cl.queue); // kernel args + int width = surface->width; - cl.err |= clSetKernelArg(contour_lines_kernel, 0, sizeof(lines_buffer), &lines_buffer); - cl.err |= clSetKernelArg(contour_lines_kernel, 1, sizeof(rows_buffer), &rows_buffer); - cl.err |= clSetKernelArg(contour_lines_kernel, 2, sizeof(mark_buffer), &mark_buffer); + cl.err |= clSetKernelArg(contour_lines_kernel, 0, sizeof(width), &width); + cl.err |= clSetKernelArg(contour_lines_kernel, 1, sizeof(lines_buffer), &lines_buffer); + cl.err |= clSetKernelArg(contour_lines_kernel, 2, sizeof(rows_buffer), &rows_buffer); + cl.err |= clSetKernelArg(contour_lines_kernel, 3, sizeof(mark_buffer), &mark_buffer); assert(!cl.err); - cl.err |= clSetKernelArg(contour_fill_kernel, 0, sizeof(mark_buffer), &mark_buffer); - cl.err |= clSetKernelArg(contour_fill_kernel, 1, sizeof(surface_buffer), &surface_buffer); - cl.err |= clSetKernelArg(contour_fill_kernel, 2, sizeof(Color::type), &color.r); - cl.err |= clSetKernelArg(contour_fill_kernel, 3, sizeof(Color::type), &color.g); - cl.err |= clSetKernelArg(contour_fill_kernel, 4, sizeof(Color::type), &color.b); - cl.err |= clSetKernelArg(contour_fill_kernel, 5, sizeof(Color::type), &color.a); + // TODO: invert, evenodd + cl.err |= clSetKernelArg(contour_fill_kernel, 0, sizeof(width), &width); + cl.err |= clSetKernelArg(contour_fill_kernel, 1, sizeof(mark_buffer), &mark_buffer); + cl.err |= clSetKernelArg(contour_fill_kernel, 2, sizeof(surface_buffer), &surface_buffer); + cl.err |= clSetKernelArg(contour_fill_kernel, 3, sizeof(Color::type), &color.r); + cl.err |= clSetKernelArg(contour_fill_kernel, 4, sizeof(Color::type), &color.g); + cl.err |= clSetKernelArg(contour_fill_kernel, 5, sizeof(Color::type), &color.b); + cl.err |= clSetKernelArg(contour_fill_kernel, 6, sizeof(Color::type), &color.a); assert(!cl.err); // prepare buffers diff --git a/c++/contourgl/clrender.h b/c++/contourgl/clrender.h index 3dd269d..874f757 100644 --- a/c++/contourgl/clrender.h +++ b/c++/contourgl/clrender.h @@ -51,7 +51,7 @@ public: void send_surface(Surface *surface); Surface* receive_surface(); - void contour(const Contour &contour, const Rect &rect, const Color &color); + void contour(const Contour &contour, const Rect &rect, const Color &color, bool invert, bool evenodd); }; diff --git a/c++/contourgl/contour.cpp b/c++/contourgl/contour.cpp index 8a5bacf..54b0866 100644 --- a/c++/contourgl/contour.cpp +++ b/c++/contourgl/contour.cpp @@ -82,8 +82,23 @@ void Contour::line_split( Rect &ref_line_bounds, const Rect &bounds, const Vector &min_size, - const Vector &p1 ) + const Vector &p1, + int level ) { + assert(level > 0); + + if (allow_split_lines) { + const Vector &p0 = current(); + if ( fabs(p1.x - p0.x) > min_size.x + || fabs(p1.y - p0.y) > min_size.y ) + { + Vector p = (p0 + p1)*0.5; + line_split(ref_line_bounds, bounds, min_size, p, level-1); + line_split(ref_line_bounds, bounds, min_size, p, level-1); + return; + } + } + line_to(p1); return; diff --git a/c++/contourgl/contour.h b/c++/contourgl/contour.h index 71f9adb..42c5c18 100644 --- a/c++/contourgl/contour.h +++ b/c++/contourgl/contour.h @@ -50,7 +50,9 @@ private: size_t first; public: - Contour(): first(0) { } + bool allow_split_lines; + + Contour(): first(0), allow_split_lines() { } void clear(); void move_to(const Vector &v); @@ -75,7 +77,8 @@ private: Rect &ref_line_bounds, const Rect &bounds, const Vector &min_size, - const Vector &p1 ); + const Vector &p1, + int level = 64 ); void conic_split( Rect &ref_line_bounds, diff --git a/c++/contourgl/measure.cpp b/c++/contourgl/measure.cpp index 06db589..cda2e07 100644 --- a/c++/contourgl/measure.cpp +++ b/c++/contourgl/measure.cpp @@ -47,6 +47,7 @@ Measure::Measure(const std::string &filename, Surface &surface): sub_tasks(), t() { + cout << string(stack.size()*2, ' ') << "begin " << filename; stack.push_back(this); t = clock(); } @@ -58,7 +59,7 @@ Measure::~Measure() { Real ms = 1000.0*(Real)(clock() - t)/(Real)(CLOCKS_PER_SEC); cout << setw(8) << fixed << setprecision(3) - << string(stack.size()*2, ' ') << "|" + << string(stack.size()*2, ' ') << "end " << ms << " ms - " << filename << endl << flush; diff --git a/c++/contourgl/test.cpp b/c++/contourgl/test.cpp index 539e368..e703cac 100644 --- a/c++/contourgl/test.cpp +++ b/c++/contourgl/test.cpp @@ -23,6 +23,7 @@ #include "triangulator.h" #include "measure.h" #include "utils.h" +#include "clrender.h" using namespace std; @@ -534,4 +535,19 @@ void Test::test4() { SwRender::polyspan(surface, polyspans[i], contours_sw[i].color, contours_sw[i].evenodd, contours_sw[i].invert); } } + + { + // cl + + vector contours_cl = contours; + Surface surface(width+2, height+2); + + Measure t("test_4_cl.tga", surface); + + ClRender clr(e.cl); + clr.send_surface(&surface); + for(vector::const_iterator i = contours_cl.begin(); i != contours_cl.end(); ++i) + clr.contour(i->contour, bounds_file, i->color, i->invert, i->evenodd); + clr.receive_surface(); + } }