Blame c++/contourgl/clrender.cpp

Ivan Mahonin 572d9c
/*
Ivan Mahonin 572d9c
    ......... 2015 Ivan Mahonin
Ivan Mahonin 572d9c
Ivan Mahonin 572d9c
    This program is free software: you can redistribute it and/or modify
Ivan Mahonin 572d9c
    it under the terms of the GNU General Public License as published by
Ivan Mahonin 572d9c
    the Free Software Foundation, either version 3 of the License, or
Ivan Mahonin 572d9c
    (at your option) any later version.
Ivan Mahonin 572d9c
Ivan Mahonin 572d9c
    This program is distributed in the hope that it will be useful,
Ivan Mahonin 572d9c
    but WITHOUT ANY WARRANTY; without even the implied warranty of
Ivan Mahonin 572d9c
    MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
Ivan Mahonin 572d9c
    GNU General Public License for more details.
Ivan Mahonin 572d9c
Ivan Mahonin 572d9c
    You should have received a copy of the GNU General Public License
Ivan Mahonin 572d9c
    along with this program.  If not, see <http://www.gnu.org/licenses/>.
Ivan Mahonin 572d9c
*/
Ivan Mahonin 572d9c
Ivan Mahonin f83e6b
#include <cassert>
Ivan Mahonin f83e6b
Ivan Mahonin f83e6b
#include <algorithm>
Ivan Mahonin f83e6b
Ivan Mahonin 572d9c
#include "clrender.h"
Ivan Mahonin f83e6b
#include "measure.h"
Ivan Mahonin 572d9c
Ivan Mahonin 572d9c
Ivan Mahonin 572d9c
using namespace std;
Ivan Mahonin 572d9c
Ivan Mahonin 572d9c
Ivan Mahonin f83e6b
ClRender::ClRender(ClContext &cl):
Ivan Mahonin f83e6b
	cl(cl),
Ivan Mahonin f83e6b
	contour_program(),
Ivan Mahonin c7fa36
	contour_path_kernel(),
Ivan Mahonin f83e6b
	contour_fill_kernel(),
Ivan Mahonin f83e6b
	surface(),
Ivan Mahonin c7fa36
	path_buffer(),
Ivan Mahonin f83e6b
	mark_buffer(),
Ivan Mahonin 67876c
	surface_image(),
Ivan Mahonin c7fa36
	prev_event()
Ivan Mahonin f83e6b
{
Ivan Mahonin f83e6b
	contour_program = cl.load_program("contour.cl");
Ivan Mahonin c7fa36
	contour_clear_kernel = clCreateKernel(contour_program, "clear", NULL);
Ivan Mahonin c7fa36
	assert(contour_clear_kernel);
Ivan Mahonin c7fa36
	contour_path_kernel = clCreateKernel(contour_program, "path", NULL);
Ivan Mahonin c7fa36
	assert(contour_path_kernel);
Ivan Mahonin f83e6b
	contour_fill_kernel = clCreateKernel(contour_program, "fill", NULL);
Ivan Mahonin f83e6b
	assert(contour_fill_kernel);
Ivan Mahonin 572d9c
}
Ivan Mahonin 572d9c
Ivan Mahonin 572d9c
ClRender::~ClRender() {
Ivan Mahonin f83e6b
	send_surface(NULL);
Ivan Mahonin c7fa36
	send_path(NULL, 0);
Ivan Mahonin c7fa36
	clReleaseKernel(contour_clear_kernel);
Ivan Mahonin f83e6b
	clReleaseKernel(contour_fill_kernel);
Ivan Mahonin c7fa36
	clReleaseKernel(contour_path_kernel);
Ivan Mahonin f83e6b
	clReleaseProgram(contour_program);
Ivan Mahonin f83e6b
}
Ivan Mahonin f83e6b
Ivan Mahonin f83e6b
void ClRender::send_surface(Surface *surface) {
Ivan Mahonin c7fa36
	if (!surface && !this->surface) return;
Ivan Mahonin f83e6b
Ivan Mahonin c7fa36
	cl.err |= clFinish(cl.queue);
Ivan Mahonin 5890eb
	assert(!cl.err);
Ivan Mahonin c7fa36
	prev_event = NULL;
Ivan Mahonin f83e6b
Ivan Mahonin f83e6b
	if (this->surface) {
Ivan Mahonin f83e6b
		clReleaseMemObject(mark_buffer);
Ivan Mahonin 67876c
		clReleaseMemObject(surface_image);
Ivan Mahonin f83e6b
	}
Ivan Mahonin f83e6b
Ivan Mahonin f83e6b
	this->surface = surface;
Ivan Mahonin f83e6b
Ivan Mahonin f83e6b
	if (this->surface) {
Ivan Mahonin 5890eb
		//Measure t("ClRender::send_surface");
Ivan Mahonin f83e6b
Ivan Mahonin 0b3288
		int width = surface->width;
Ivan Mahonin 0b3288
		int height = surface->height;
Ivan Mahonin 0b3288
Ivan Mahonin f83e6b
		mark_buffer = clCreateBuffer(
Ivan Mahonin f83e6b
			cl.context, CL_MEM_READ_WRITE,
Ivan Mahonin 0b3288
			surface->count()*sizeof(cl_int4), NULL,
Ivan Mahonin f83e6b
			NULL );
Ivan Mahonin f83e6b
		assert(mark_buffer);
Ivan Mahonin f83e6b
Ivan Mahonin c7fa36
		cl.err |= clSetKernelArg(contour_clear_kernel, 0, sizeof(mark_buffer), &mark_buffer);
Ivan Mahonin 0b3288
		cl.err |= clSetKernelArg(contour_clear_kernel, 1, sizeof(width), &width);
Ivan Mahonin c7fa36
		assert(!cl.err);
Ivan Mahonin c7fa36
Ivan Mahonin c7fa36
		size_t pixels_count = (size_t)surface->count();
Ivan Mahonin c7fa36
		cl.err |= clEnqueueNDRangeKernel(
Ivan Mahonin c7fa36
			cl.queue,
Ivan Mahonin c7fa36
			contour_clear_kernel,
Ivan Mahonin c7fa36
			1,
Ivan Mahonin c7fa36
			NULL,
Ivan Mahonin c7fa36
			&pixels_count,
Ivan Mahonin c7fa36
			NULL,
Ivan Mahonin c7fa36
			0,
Ivan Mahonin c7fa36
			NULL,
Ivan Mahonin c7fa36
			NULL );
Ivan Mahonin c7fa36
		assert(!cl.err);
Ivan Mahonin c7fa36
Ivan Mahonin 67876c
		cl_image_format surface_format = { };
Ivan Mahonin 67876c
		surface_format.image_channel_order = CL_RGBA;
Ivan Mahonin 67876c
		surface_format.image_channel_data_type = CL_FLOAT;
Ivan Mahonin f83e6b
Ivan Mahonin 67876c
		surface_image = clCreateImage2D(
Ivan Mahonin 67876c
			cl.context, CL_MEM_READ_WRITE,
Ivan Mahonin 67876c
			&surface_format, surface->width, surface->height,
Ivan Mahonin 67876c
			0, NULL, NULL );
Ivan Mahonin 67876c
		assert(surface_image);
Ivan Mahonin 67876c
Ivan Mahonin 67876c
		size_t origin[3] = { };
Ivan Mahonin 67876c
		size_t region[3] = { (size_t)surface->width, (size_t)surface->height, 1 };
Ivan Mahonin 67876c
		cl.err |= clEnqueueWriteImage(
Ivan Mahonin c7fa36
			cl.queue, surface_image, CL_FALSE,
Ivan Mahonin 67876c
			origin, region, 0, 0, surface->data,
Ivan Mahonin c7fa36
			0, NULL, NULL );
Ivan Mahonin c7fa36
		assert(!cl.err);
Ivan Mahonin c7fa36
Ivan Mahonin c7fa36
		cl.err |= clSetKernelArg(contour_path_kernel, 0, sizeof(width), &width);
Ivan Mahonin c7fa36
		cl.err |= clSetKernelArg(contour_path_kernel, 1, sizeof(height), &height);
Ivan Mahonin c7fa36
		cl.err |= clSetKernelArg(contour_path_kernel, 2, sizeof(mark_buffer), &mark_buffer);
Ivan Mahonin c7fa36
		assert(!cl.err);
Ivan Mahonin c7fa36
Ivan Mahonin c7fa36
		cl.err |= clSetKernelArg(contour_fill_kernel, 0, sizeof(width), &width);
Ivan Mahonin c7fa36
		cl.err |= clSetKernelArg(contour_fill_kernel, 1, sizeof(mark_buffer), &mark_buffer);
Ivan Mahonin c7fa36
		cl.err |= clSetKernelArg(contour_fill_kernel, 2, sizeof(surface_image), &surface_image);
Ivan Mahonin c7fa36
		cl.err |= clSetKernelArg(contour_fill_kernel, 3, sizeof(surface_image), &surface_image);
Ivan Mahonin c7fa36
		assert(!cl.err);
Ivan Mahonin f83e6b
Ivan Mahonin c7fa36
		cl.err |= clFinish(cl.queue);
Ivan Mahonin 5890eb
		assert(!cl.err);
Ivan Mahonin f83e6b
	}
Ivan Mahonin f83e6b
}
Ivan Mahonin f83e6b
Ivan Mahonin f83e6b
Surface* ClRender::receive_surface() {
Ivan Mahonin f83e6b
	if (surface) {
Ivan Mahonin 5890eb
		//Measure t("ClRender::receive_surface");
Ivan Mahonin f83e6b
Ivan Mahonin 67876c
		size_t origin[3] = { };
Ivan Mahonin 67876c
		size_t region[3] = { (size_t)surface->width, (size_t)surface->height, 1 };
Ivan Mahonin 67876c
		cl.err |= clEnqueueReadImage(
Ivan Mahonin c7fa36
			cl.queue, surface_image, CL_FALSE,
Ivan Mahonin 67876c
			origin, region, 0, 0, surface->data,
Ivan Mahonin 028154
			prev_event ? 1 : 0, &prev_event, NULL );
Ivan Mahonin 5890eb
		assert(!cl.err);
Ivan Mahonin c7fa36
Ivan Mahonin c7fa36
		cl.err |= clFinish(cl.queue);
Ivan Mahonin c7fa36
		assert(!cl.err);
Ivan Mahonin 028154
		prev_event = NULL;
Ivan Mahonin f83e6b
	}
Ivan Mahonin f83e6b
	return surface;
Ivan Mahonin 572d9c
}
Ivan Mahonin 572d9c
Ivan Mahonin c7fa36
void ClRender::send_path(const vec2f *path, int count) {
Ivan Mahonin c7fa36
	if (!path_buffer && (!path || count <=0)) return;
Ivan Mahonin f83e6b
Ivan Mahonin c7fa36
	cl.err |= clFinish(cl.queue);
Ivan Mahonin c7fa36
	assert(!cl.err);
Ivan Mahonin c7fa36
	prev_event = NULL;
Ivan Mahonin f83e6b
Ivan Mahonin c7fa36
	if (path_buffer) {
Ivan Mahonin c7fa36
		clReleaseMemObject(path_buffer);
Ivan Mahonin c7fa36
		path_buffer = NULL;
Ivan Mahonin f83e6b
	}
Ivan Mahonin f83e6b
Ivan Mahonin c7fa36
	if (path && count > 0) {
Ivan Mahonin c7fa36
		//Measure t("ClRender::send_path");
Ivan Mahonin f83e6b
Ivan Mahonin c7fa36
		path_buffer = clCreateBuffer(
Ivan Mahonin f83e6b
			cl.context, CL_MEM_READ_ONLY,
Ivan Mahonin c7fa36
			count*sizeof(*path), NULL,
Ivan Mahonin f83e6b
			NULL );
Ivan Mahonin c7fa36
		assert(path_buffer);
Ivan Mahonin f83e6b
Ivan Mahonin f83e6b
		cl.err |= clEnqueueWriteBuffer(
Ivan Mahonin c7fa36
			cl.queue, path_buffer, CL_FALSE,
Ivan Mahonin c7fa36
			0, count*sizeof(*path), path,
Ivan Mahonin c7fa36
			0, NULL, NULL );
Ivan Mahonin 5890eb
		assert(!cl.err);
Ivan Mahonin f83e6b
Ivan Mahonin c7fa36
		cl.err |= clSetKernelArg(contour_path_kernel, 3, sizeof(path_buffer), &path_buffer);
Ivan Mahonin 5890eb
		assert(!cl.err);
Ivan Mahonin f83e6b
Ivan Mahonin c7fa36
		cl.err |= clFinish(cl.queue);
Ivan Mahonin 5890eb
		assert(!cl.err);
Ivan Mahonin c7fa36
	}
Ivan Mahonin c7fa36
}
Ivan Mahonin f83e6b
Ivan Mahonin c7fa36
void ClRender::path(int start, int count, const Color &color, bool invert, bool evenodd) {
Ivan Mahonin c7fa36
	//Measure t("ClRender::contour");
Ivan Mahonin f83e6b
Ivan Mahonin c7fa36
	if (count <= 1) return;
Ivan Mahonin f83e6b
Ivan Mahonin c7fa36
	// kernel args
Ivan Mahonin 0b3288
Ivan Mahonin c7fa36
	int iinvert = invert, ievenodd = evenodd;
Ivan Mahonin c7fa36
	cl.err |= clSetKernelArg(contour_fill_kernel, 4, sizeof(color), &color);
Ivan Mahonin c7fa36
	cl.err |= clSetKernelArg(contour_fill_kernel, 5, sizeof(int), &iinvert);
Ivan Mahonin c7fa36
	cl.err |= clSetKernelArg(contour_fill_kernel, 6, sizeof(int), &ievenodd);
Ivan Mahonin c7fa36
	assert(!cl.err);
Ivan Mahonin 67876c
Ivan Mahonin c7fa36
	// build marks
Ivan Mahonin c7fa36
Ivan Mahonin c7fa36
	cl_event path_event = NULL;
Ivan Mahonin c7fa36
	size_t sstart = start;
Ivan Mahonin c7fa36
	size_t scount = count-1;
Ivan Mahonin c7fa36
	cl.err |= clEnqueueNDRangeKernel(
Ivan Mahonin c7fa36
		cl.queue,
Ivan Mahonin c7fa36
		contour_path_kernel,
Ivan Mahonin c7fa36
		1,
Ivan Mahonin c7fa36
		&sstart,
Ivan Mahonin c7fa36
		&scount,
Ivan Mahonin c7fa36
		NULL,
Ivan Mahonin c7fa36
		prev_event ? 1 : 0,
Ivan Mahonin c7fa36
		&prev_event,
Ivan Mahonin c7fa36
		&path_event );
Ivan Mahonin c7fa36
	assert(!cl.err);
Ivan Mahonin f83e6b
Ivan Mahonin c7fa36
	// fill
Ivan Mahonin c7fa36
	size_t sheight = surface->height;
Ivan Mahonin c7fa36
	cl.err |= clEnqueueNDRangeKernel(
Ivan Mahonin c7fa36
		cl.queue,
Ivan Mahonin c7fa36
		contour_fill_kernel,
Ivan Mahonin c7fa36
		1,
Ivan Mahonin c7fa36
		NULL,
Ivan Mahonin c7fa36
		&sheight,
Ivan Mahonin c7fa36
		NULL,
Ivan Mahonin c7fa36
		1,
Ivan Mahonin c7fa36
		&path_event,
Ivan Mahonin c7fa36
		&prev_event );
Ivan Mahonin c7fa36
	assert(!cl.err);
Ivan Mahonin 572d9c
}
Ivan Mahonin 572d9c
Ivan Mahonin 572d9c
Ivan Mahonin 572d9c
Ivan Mahonin 572d9c
void SwRenderAlt::line(const Vector &p0, const Vector &p1) {
Ivan Mahonin 572d9c
	int iy0 = min(max((int)floor(p0.y), 0), height);
Ivan Mahonin 572d9c
	int iy1 = min(max((int)floor(p1.y), 0), height);
Ivan Mahonin 572d9c
	if (iy1 < iy0) swap(iy0, iy1);
Ivan Mahonin 572d9c
Ivan Mahonin 572d9c
	Vector d = p1 - p0;
Ivan Mahonin 572d9c
	Vector k( fabs(d.y) < 1e-6 ? 0.0 : d.x/d.y,
Ivan Mahonin 572d9c
		      fabs(d.x) < 1e-6 ? 0.0 : d.y/d.x );
Ivan Mahonin 572d9c
Ivan Mahonin 572d9c
	for(int r = iy0; r <= iy1; ++r) {
Ivan Mahonin 572d9c
		Real y = (Real)iy0;
Ivan Mahonin 572d9c
Ivan Mahonin 572d9c
		Vector pp0 = p0;
Ivan Mahonin 572d9c
		pp0.y -= y;
Ivan Mahonin 572d9c
		if (pp0.y < 0.0) {
Ivan Mahonin 572d9c
			pp0.y = 0.0;
Ivan Mahonin 572d9c
			pp0.x = p0.x - k.x*y;
Ivan Mahonin 572d9c
		} else
Ivan Mahonin 572d9c
		if (pp0.y > 1.0) {
Ivan Mahonin 572d9c
			pp0.y = 1.0;
Ivan Mahonin 572d9c
			pp0.x = p0.x - k.x*(y - 1.0);
Ivan Mahonin 572d9c
		}
Ivan Mahonin 572d9c
Ivan Mahonin 572d9c
		Vector pp1 = p1;
Ivan Mahonin 572d9c
		pp1.y -= y;
Ivan Mahonin 572d9c
		if (pp1.y < 0.0) {
Ivan Mahonin 572d9c
			pp1.y = 0.0;
Ivan Mahonin 572d9c
			pp1.x = p0.x - k.x*y;
Ivan Mahonin 572d9c
		} else
Ivan Mahonin 572d9c
		if (pp1.y > 1.0) {
Ivan Mahonin 572d9c
			pp1.y = 1.0;
Ivan Mahonin 572d9c
			pp1.x = p0.x - k.x*(y - 1.0);
Ivan Mahonin 572d9c
		}
Ivan Mahonin 572d9c
Ivan Mahonin 572d9c
		int ix0 = min(max((int)floor(pp0.x), 0), width);
Ivan Mahonin 572d9c
		int ix1 = min(max((int)floor(pp1.x), 0), width);
Ivan Mahonin 572d9c
		if (ix1 < ix0) swap(ix0, ix1);
Ivan Mahonin 572d9c
		for(int c = ix0; c <= ix1; ++c) {
Ivan Mahonin 572d9c
			Real x = (Real)ix0;
Ivan Mahonin 572d9c
Ivan Mahonin 572d9c
			Vector ppp0 = pp0;
Ivan Mahonin 572d9c
			ppp0.x -= x;
Ivan Mahonin 572d9c
			if (ppp0.x < 0.0) {
Ivan Mahonin 572d9c
				ppp0.x = 0.0;
Ivan Mahonin 572d9c
				ppp0.y = pp0.y - k.y*x;
Ivan Mahonin 572d9c
			} else
Ivan Mahonin 572d9c
			if (ppp0.x > 1.0) {
Ivan Mahonin 572d9c
				ppp0.x = 1.0;
Ivan Mahonin 572d9c
				ppp0.y = pp0.y - k.y*(x - 1.0);
Ivan Mahonin 572d9c
			}
Ivan Mahonin 572d9c
Ivan Mahonin 572d9c
			Vector ppp1 = pp1;
Ivan Mahonin 572d9c
			ppp1.x -= x;
Ivan Mahonin 572d9c
			if (ppp1.x < 0.0) {
Ivan Mahonin 572d9c
				ppp1.x = 0.0;
Ivan Mahonin 572d9c
				ppp1.y = pp0.y - k.y*x;
Ivan Mahonin 572d9c
			} else
Ivan Mahonin 572d9c
			if (ppp1.x > 1.0) {
Ivan Mahonin 572d9c
				ppp1.x = 1.0;
Ivan Mahonin 572d9c
				ppp1.y = pp0.y - k.y*(x - 1.0);
Ivan Mahonin 572d9c
			}
Ivan Mahonin 572d9c
Ivan Mahonin 572d9c
			Real cover = ppp0.y - ppp1.y;
Ivan Mahonin 572d9c
			Real area = (0.5*(ppp1.x + ppp1.x) - 1.0)*cover;
Ivan Mahonin 572d9c
			(*this)[r][c].add(area, cover);
Ivan Mahonin 572d9c
		}
Ivan Mahonin 572d9c
	}
Ivan Mahonin 572d9c
}