/*

    ......... 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 <http: licenses="" www.gnu.org="">.</http:>

*/

#include <cassert></cassert>

#include <algorithm></algorithm>

#include "clrender.h"

#include "measure.h"

using namespace std;

ClRender::ClRender(ClContext &cl):

	cl(cl),

	contour_program(),

	contour_path_kernel(),

	contour_fill_kernel(),

	surface(),

	path_buffer(),

	mark_buffer(),

	surface_image(),

	prev_event()

{

	contour_program = cl.load_program("contour.cl");

	contour_clear_kernel = clCreateKernel(contour_program, "clear", NULL);

	assert(contour_clear_kernel);

	contour_path_kernel = clCreateKernel(contour_program, "path", NULL);

	assert(contour_path_kernel);

	contour_fill_kernel = clCreateKernel(contour_program, "fill", NULL);

	assert(contour_fill_kernel);

}

ClRender::~ClRender() {

	send_surface(NULL);

	send_path(NULL, 0);

	clReleaseKernel(contour_clear_kernel);

	clReleaseKernel(contour_fill_kernel);

	clReleaseKernel(contour_path_kernel);

	clReleaseProgram(contour_program);

}

void ClRender::send_surface(Surface *surface) {

	if (!surface && !this->surface) return;

	cl.err |= clFinish(cl.queue);

	assert(!cl.err);

	prev_event = NULL;

	if (this->surface) {

		clReleaseMemObject(mark_buffer);

		clReleaseMemObject(surface_image);

	}

	this->surface = surface;

	if (this->surface) {

		//Measure t("ClRender::send_surface");

		int width = surface->width;

		int height = surface->height;

		mark_buffer = clCreateBuffer(

			cl.context, CL_MEM_READ_WRITE,

			surface->count()*sizeof(cl_int4), NULL,

			NULL );

		assert(mark_buffer);

		cl.err |= clSetKernelArg(contour_clear_kernel, 0, sizeof(mark_buffer), &mark_buffer);

		cl.err |= clSetKernelArg(contour_clear_kernel, 1, sizeof(width), &width);

		assert(!cl.err);

		size_t pixels_count = (size_t)surface->count();

		cl.err |= clEnqueueNDRangeKernel(

			cl.queue,

			contour_clear_kernel,

			1,

			NULL,

			&pixels_count,

			NULL,

			0,

			NULL,

			NULL );

		assert(!cl.err);

		cl_image_format surface_format = { };

		surface_format.image_channel_order = CL_RGBA;

		surface_format.image_channel_data_type = CL_FLOAT;

		surface_image = clCreateImage2D(

			cl.context, CL_MEM_READ_WRITE,

			&surface_format, surface->width, surface->height,

			0, NULL, NULL );

		assert(surface_image);

		size_t origin[3] = { };

		size_t region[3] = { (size_t)surface->width, (size_t)surface->height, 1 };

		cl.err |= clEnqueueWriteImage(

			cl.queue, surface_image, CL_FALSE,

			origin, region, 0, 0, surface->data,

			0, NULL, NULL );

		assert(!cl.err);

		cl.err |= clSetKernelArg(contour_path_kernel, 0, sizeof(width), &width);

		cl.err |= clSetKernelArg(contour_path_kernel, 1, sizeof(height), &height);

		cl.err |= clSetKernelArg(contour_path_kernel, 2, sizeof(mark_buffer), &mark_buffer);

		assert(!cl.err);

		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_image), &surface_image);

		cl.err |= clSetKernelArg(contour_fill_kernel, 3, sizeof(surface_image), &surface_image);

		assert(!cl.err);

		cl.err |= clFinish(cl.queue);

		assert(!cl.err);

	}

}

Surface* ClRender::receive_surface() {

	if (surface) {

		//Measure t("ClRender::receive_surface");

		size_t origin[3] = { };

		size_t region[3] = { (size_t)surface->width, (size_t)surface->height, 1 };

		cl.err |= clEnqueueReadImage(

			cl.queue, surface_image, CL_FALSE,

			origin, region, 0, 0, surface->data,

			prev_event ? 1 : 0, &prev_event, NULL );

		assert(!cl.err);

		cl.err |= clFinish(cl.queue);

		assert(!cl.err);

		prev_event = NULL;

	}

	return surface;

}

void ClRender::send_path(const vec2f *path, int count) {

	if (!path_buffer && (!path || count <=0)) return;

	cl.err |= clFinish(cl.queue);

	assert(!cl.err);

	prev_event = NULL;

	if (path_buffer) {

		clReleaseMemObject(path_buffer);

		path_buffer = NULL;

	}

	if (path && count > 0) {

		//Measure t("ClRender::send_path");

		path_buffer = clCreateBuffer(

			cl.context, CL_MEM_READ_ONLY,

			count*sizeof(*path), NULL,

			NULL );

		assert(path_buffer);

		cl.err |= clEnqueueWriteBuffer(

			cl.queue, path_buffer, CL_FALSE,

			0, count*sizeof(*path), path,

			0, NULL, NULL );

		assert(!cl.err);

		cl.err |= clSetKernelArg(contour_path_kernel, 3, sizeof(path_buffer), &path_buffer);

		assert(!cl.err);

		cl.err |= clFinish(cl.queue);

		assert(!cl.err);

	}

}

void ClRender::path(int start, int count, const Color &color, bool invert, bool evenodd) {

	//Measure t("ClRender::contour");

	if (count <= 1) return;

	// kernel args

	int iinvert = invert, ievenodd = evenodd;

	cl.err |= clSetKernelArg(contour_fill_kernel, 4, sizeof(color), &color);

	cl.err |= clSetKernelArg(contour_fill_kernel, 5, sizeof(int), &iinvert);

	cl.err |= clSetKernelArg(contour_fill_kernel, 6, sizeof(int), &ievenodd);

	assert(!cl.err);

	// build marks

	cl_event path_event = NULL;

	size_t sstart = start;

	size_t scount = count-1;

	cl.err |= clEnqueueNDRangeKernel(

		cl.queue,

		contour_path_kernel,

		1,

		&sstart,

		&scount,

		NULL,

		prev_event ? 1 : 0,

		&prev_event,

		&path_event );

	assert(!cl.err);

	// fill

	size_t sheight = surface->height;

	cl.err |= clEnqueueNDRangeKernel(

		cl.queue,

		contour_fill_kernel,

		1,

		NULL,

		&sheight,

		NULL,

		1,

		&path_event,

		&prev_event );

	assert(!cl.err);

}

void ClRender::wait() {

	if (prev_event) {

		clWaitForEvents(1, &prev_event);

		prev_event = NULL;

	}

}

void SwRenderAlt::line(const Vector &p0, const Vector &p1) {

	int iy0 = min(max((int)floor(p0.y), 0), height);

	int iy1 = min(max((int)floor(p1.y), 0), height);

	if (iy1 < iy0) swap(iy0, iy1);

	Vector d = p1 - p0;

	Vector 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) {

		Real y = (Real)iy0;

		Vector 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);

		}

		Vector 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) swap(ix0, ix1);

		for(int c = ix0; c <= ix1; ++c) {

			Real x = (Real)ix0;

			Vector 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);

			}

			Vector 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);

			}

			Real cover = ppp0.y - ppp1.y;

			Real area = (0.5*(ppp1.x + ppp1.x) - 1.0)*cover;

			(*this)[r][c].add(area, cover);

		}

	}

}