/*
......... 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://www.gnu.org/licenses/>.
*/
#include <cassert>
#include <algorithm>
#include <iostream>
#include "clrender.h"
#include "measure.h"
using namespace std;
ClRender::ClRender(ClContext &cl):
cl(cl),
contour_program(),
contour_draw_kernel(),
contour_draw_workgroup_size(),
surface(),
paths_buffer(),
mark_buffer(),
surface_image(),
prev_event()
{
contour_program = cl.load_program("contour-fs.cl");
assert(contour_program);
contour_draw_kernel = clCreateKernel(contour_program, "draw", NULL);
assert(contour_draw_kernel);
cl.err |= clGetKernelWorkGroupInfo(
contour_draw_kernel,
cl.device,
CL_KERNEL_WORK_GROUP_SIZE,
sizeof(contour_draw_workgroup_size),
&contour_draw_workgroup_size,
NULL );
assert(!cl.err);
}
ClRender::~ClRender() {
send_surface(NULL);
send_paths(NULL, 0);
clReleaseKernel(contour_draw_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");
mark_buffer = clCreateBuffer(
cl.context, CL_MEM_READ_WRITE,
(surface->count() + 2)*sizeof(cl_int2), NULL,
&cl.err );
assert(!cl.err);
assert(mark_buffer);
char zero = 0;
cl.err |= clEnqueueFillBuffer(
cl.queue, mark_buffer,
&zero, 1,
0, surface->count()*sizeof(cl_int2),
0, NULL, NULL );
assert(!cl.err);
surface_image = clCreateBuffer(
cl.context, CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR,
surface->count()*sizeof(Color), surface->data,
&cl.err );
assert(!cl.err);
assert(surface_image);
cl.err |= clSetKernelArg(contour_draw_kernel, 0, sizeof(surface->width), &surface->width);
cl.err |= clSetKernelArg(contour_draw_kernel, 1, sizeof(surface->width), &surface->height);
cl.err |= clSetKernelArg(contour_draw_kernel, 2, sizeof(mark_buffer), &mark_buffer);
cl.err |= clSetKernelArg(contour_draw_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");
cl.err |= clEnqueueReadBuffer(
cl.queue, surface_image, CL_FALSE,
0, surface->count()*sizeof(Color), surface->data,
prev_event ? 1 : 0,
prev_event ? &prev_event : NULL,
NULL );
assert(!cl.err);
cl.err |= clFinish(cl.queue);
assert(!cl.err);
prev_event = NULL;
}
return surface;
}
void ClRender::remove_paths() {
if (paths_buffer) {
cl.err |= clFinish(cl.queue);
assert(!cl.err);
prev_event = NULL;
clReleaseMemObject(paths_buffer);
paths_buffer = NULL;
}
}
void ClRender::send_paths(const void *paths, int size) {
if (!paths_buffer && (!paths || size <= 0)) return;
remove_paths();
if (paths && size > 0) {
//Measure t("ClRender::send_path");
paths_buffer = clCreateBuffer(
cl.context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,
size, const_cast<void*>(paths),
&cl.err );
assert(!cl.err);
assert(paths_buffer);
cl.err |= clSetKernelArg(contour_draw_kernel, 4, sizeof(paths_buffer), &paths_buffer);
assert(!cl.err);
}
}
void ClRender::draw() {
//Measure t("ClRender::contour");
cl_event event = prev_event;
size_t count = contour_draw_workgroup_size;
size_t group_size = count;
cl.err |= clEnqueueNDRangeKernel(
cl.queue,
contour_draw_kernel,
1,
NULL,
&count,
&group_size,
event ? 1 : 0,
event ? &event : NULL,
&prev_event );
assert(!cl.err);
}
void ClRender::wait() {
if (prev_event) {
cl.err |= clWaitForEvents(1, &prev_event);
assert(!cl.err);
prev_event = NULL;
}
}
// ------------------------------------------------
ClRender2::ClRender2(ClContext &cl):
cl(cl),
contour_program(),
contour_reset_kernel(),
contour_paths_kernel(),
contour_draw_kernel(),
surface(),
points_count(),
paths_buffer(),
points_buffer(),
samples_buffer(),
surface_image(),
prev_event()
{
contour_program = cl.load_program("contour-sort.cl");
assert(contour_program);
contour_reset_kernel = clCreateKernel(contour_program, "reset", &cl.err);
assert(!cl.err);
assert(contour_reset_kernel);
contour_paths_kernel = clCreateKernel(contour_program, "paths", &cl.err);
assert(!cl.err);
assert(contour_paths_kernel);
contour_draw_kernel = clCreateKernel(contour_program, "draw", &cl.err);
assert(!cl.err);
assert(contour_draw_kernel);
samples_buffer = clCreateBuffer(
cl.context, CL_MEM_READ_WRITE,
1024*1024*1024, NULL,
&cl.err );
assert(!cl.err);
assert(samples_buffer);
cl.err |= clSetKernelArg(contour_reset_kernel, 0, sizeof(samples_buffer), &samples_buffer);
cl.err |= clSetKernelArg(contour_paths_kernel, 2, sizeof(samples_buffer), &samples_buffer);
cl.err |= clSetKernelArg(contour_draw_kernel, 2, sizeof(samples_buffer), &samples_buffer);
assert(!cl.err);
}
ClRender2::~ClRender2() {
remove_paths();
remove_surface();
cl.err |= clReleaseMemObject(samples_buffer);
assert(!cl.err);
samples_buffer = NULL;
clReleaseKernel(contour_reset_kernel);
clReleaseKernel(contour_paths_kernel);
clReleaseKernel(contour_draw_kernel);
clReleaseProgram(contour_program);
}
void ClRender2::remove_surface() {
wait();
if (surface) {
cl.err |= clReleaseMemObject(surface_image);
assert(!cl.err);
surface = NULL;
}
}
void ClRender2::send_surface(Surface *surface) {
if (!surface && !this->surface) return;
remove_surface();
assert(surface);
this->surface = surface;
//Measure t("ClRender::send_surface");
surface_image = clCreateBuffer(
cl.context, CL_MEM_READ_WRITE,
surface->count()*sizeof(Color), NULL,
&cl.err );
assert(!cl.err);
assert(surface_image);
cl.err |= clEnqueueWriteBuffer(
cl.queue, surface_image, false,
0, surface->count()*sizeof(Color), surface->data,
0, NULL, NULL );
assert(!cl.err);
cl.err |= clSetKernelArg(contour_paths_kernel, 0, sizeof(surface->width), &surface->width);
cl.err |= clSetKernelArg(contour_paths_kernel, 1, sizeof(surface->height), &surface->height);
cl.err |= clSetKernelArg(contour_draw_kernel, 0, sizeof(surface->width), &surface->width);
cl.err |= clSetKernelArg(contour_draw_kernel, 1, sizeof(surface_image), &surface_image);
assert(!cl.err);
}
Surface* ClRender2::receive_surface() {
if (surface) {
//Measure t("ClRender::receive_surface");
cl.err |= clEnqueueReadBuffer(
cl.queue, surface_image, CL_FALSE,
0, surface->count()*sizeof(Color), surface->data,
prev_event ? 1 : 0,
prev_event ? &prev_event : NULL,
NULL );
assert(!cl.err);
wait();
}
return surface;
}
void ClRender2::remove_paths() {
wait();
if (paths_buffer) {
cl.err |= clReleaseMemObject(paths_buffer);
assert(!cl.err);
paths_buffer = NULL;
}
if (points_buffer) {
cl.err |= clReleaseMemObject(points_buffer);
assert(!cl.err);
points_buffer = NULL;
points_count = 0;
}
}
void ClRender2::send_paths(const Path *paths, int paths_count, const Point *points, int points_count) {
remove_paths();
assert(paths);
assert(paths_count > 0);
assert(points);
assert(points_count > 0);
paths_buffer = clCreateBuffer(
cl.context, CL_MEM_READ_ONLY,
paths_count*sizeof(Path), NULL,
&cl.err );
assert(!cl.err);
assert(paths_buffer);
cl.err |= clEnqueueWriteBuffer(
cl.queue, paths_buffer, false,
0, paths_count*sizeof(Path), paths,
0, NULL, NULL );
assert(!cl.err);
points_buffer = clCreateBuffer(
cl.context, CL_MEM_READ_ONLY,
points_count*sizeof(Point), NULL,
&cl.err );
assert(!cl.err);
assert(points_buffer);
this->points_count = points_count;
cl.err |= clEnqueueWriteBuffer(
cl.queue, points_buffer, false,
0, points_count*sizeof(Point), points,
0, NULL, NULL );
assert(!cl.err);
cl.err |= clSetKernelArg(contour_paths_kernel, 3, sizeof(points_buffer), &points_buffer);
cl.err |= clSetKernelArg(contour_draw_kernel, 3, sizeof(paths_buffer), &paths_buffer);
assert(!cl.err);
wait();
}
void ClRender2::draw() {
//Measure t("ClRender::contour");
cl_event prepare_event;
cl_event paths_event;
size_t count = surface->height;
cl.err |= clEnqueueNDRangeKernel(
cl.queue,
contour_reset_kernel,
1,
NULL,
&count,
NULL,
prev_event ? 1 : 0,
prev_event ? &prev_event : NULL,
&prepare_event );
assert(!cl.err);
count = points_count - 1;
cl.err |= clEnqueueNDRangeKernel(
cl.queue,
contour_paths_kernel,
1,
NULL,
&count,
NULL,
1,
&prepare_event,
&paths_event );
assert(!cl.err);
count = surface->height;
cl.err |= clEnqueueNDRangeKernel(
cl.queue,
contour_draw_kernel,
1,
NULL,
&count,
NULL,
1,
&paths_event,
&prev_event );
assert(!cl.err);
}
void ClRender2::wait() {
cl.err |= clFinish(cl.queue);
assert(!cl.err);
prev_event = NULL;
}
// ------------------------------------------------
ClRender3::ClRender3(ClContext &cl):
cl(cl),
contour_program(),
contour_path_kernel(),
contour_fill_kernel(),
surface(),
points_buffer(),
mark_buffer(),
surface_image(),
prev_event()
{
contour_program = cl.load_program("contour-base.cl");
assert(contour_program);
contour_path_kernel = clCreateKernel(contour_program, "path", &cl.err);
assert(!cl.err);
assert(contour_path_kernel);
contour_fill_kernel = clCreateKernel(contour_program, "fill", &cl.err);
assert(!cl.err);
assert(contour_fill_kernel);
}
ClRender3::~ClRender3() {
send_points(NULL, 0);
send_surface(NULL);
cl.err |= clReleaseKernel(contour_path_kernel);
cl.err |= clReleaseKernel(contour_fill_kernel);
cl.err |= clReleaseProgram(contour_program);
assert(!cl.err);
}
void ClRender3::send_surface(Surface *surface) {
if (this->surface) {
wait();
cl.err |= clReleaseMemObject(surface_image);
assert(!cl.err);
surface_image = NULL;
}
this->surface = surface;
if (this->surface) {
//Measure t("ClRender::send_surface");
vec2i zero_mark;
surface_image = clCreateBuffer(
cl.context, CL_MEM_READ_WRITE,
surface->count()*sizeof(Color), NULL,
&cl.err );
assert(!cl.err);
assert(surface_image);
mark_buffer = clCreateBuffer(
cl.context, CL_MEM_READ_WRITE,
surface->count()*sizeof(zero_mark), NULL,
&cl.err );
assert(!cl.err);
assert(mark_buffer);
cl.err |= clEnqueueWriteBuffer(
cl.queue, surface_image, false,
0, surface->count()*sizeof(Color), surface->data,
0, NULL, NULL );
assert(!cl.err);
cl.err |= clEnqueueFillBuffer(
cl.queue, mark_buffer,
&zero_mark, sizeof(zero_mark),
0, surface->count()*sizeof(zero_mark),
0, NULL, NULL );
assert(!cl.err);
cl.err |= clSetKernelArg(contour_path_kernel, 0, sizeof(surface->width), &surface->width);
cl.err |= clSetKernelArg(contour_path_kernel, 1, sizeof(surface->height), &surface->height);
cl.err |= clSetKernelArg(contour_path_kernel, 2, sizeof(mark_buffer), &mark_buffer);
assert(!cl.err);
cl.err |= clSetKernelArg(contour_fill_kernel, 0, sizeof(surface->width), &surface->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);
assert(!cl.err);
wait();
}
}
Surface* ClRender3::receive_surface() {
if (surface) {
//Measure t("ClRender::receive_surface");
cl.err |= clEnqueueReadBuffer(
cl.queue, surface_image, CL_FALSE,
0, surface->count()*sizeof(Color), surface->data,
prev_event ? 1 : 0,
prev_event ? &prev_event : NULL,
NULL );
assert(!cl.err);
wait();
}
return surface;
}
void ClRender3::send_points(const vec2f *points, int count) {
if (points_buffer) {
wait();
cl.err |= clReleaseMemObject(points_buffer);
assert(!cl.err);
points_buffer = NULL;
}
if (points && count > 0) {
points_buffer = clCreateBuffer(
cl.context, CL_MEM_READ_ONLY,
count*sizeof(vec2f), NULL,
&cl.err );
assert(!cl.err);
assert(points_buffer);
cl.err |= clEnqueueWriteBuffer(
cl.queue, points_buffer, false,
0, count*sizeof(vec2f), points,
0, NULL, NULL );
assert(!cl.err);
cl.err |= clSetKernelArg(contour_path_kernel, 3, sizeof(points_buffer), &points_buffer);
assert(!cl.err);
wait();
}
}
void ClRender3::draw(const Path &path) {
//Measure t("ClRender::contour");
assert(surface);
assert(points_buffer);
ContextRect bounds;
bounds.minx = max(1, path.bounds.minx);
bounds.maxx = min(surface->width, path.bounds.maxx);
bounds.miny = max(0, path.bounds.miny);
bounds.maxy = min(surface->height, path.bounds.maxy);
if ( bounds.minx >= bounds.maxx
|| bounds.miny >= bounds.maxy
|| path.begin >= path.end ) return;
vec2i boundsx(bounds.minx, bounds.maxx);
cl.err |= clSetKernelArg(contour_path_kernel, 4, sizeof(path.end), &path.end);
cl.err |= clSetKernelArg(contour_path_kernel, 5, sizeof(bounds.minx), &bounds.minx);
assert(!cl.err);
cl.err |= clSetKernelArg(contour_fill_kernel, 3, sizeof(path.color), &path.color);
cl.err |= clSetKernelArg(contour_fill_kernel, 4, sizeof(bounds), &bounds);
assert(!cl.err);
size_t group_size, offset, count;
offset = path.begin;
count = path.end - path.begin - 1;
group_size = 128;
count = ((count - 1)/group_size + 1)*group_size;
cl.err |= clEnqueueNDRangeKernel(
cl.queue, contour_path_kernel,
1, &offset, &count, &group_size,
0, NULL, NULL );
assert(!cl.err);
offset = bounds.minx;
count = bounds.maxx - bounds.minx;
group_size = 16;
count = ((count - 1)/group_size + 1)*group_size;
cl.err |= clEnqueueNDRangeKernel(
cl.queue, contour_fill_kernel,
1, &offset, &count, &group_size,
0, NULL, NULL );
assert(!cl.err);
}
void ClRender3::wait() {
cl.err |= clFinish(cl.queue);
assert(!cl.err);
if (prev_event) {
cl.err |= clReleaseEvent(prev_event);
assert(!cl.err);
prev_event = NULL;
}
}