Blame c++/contourgl/clrender.cpp

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