Blame c++/contourgl/cudarender.cpp

6fa009
/*
6fa009
    ......... 2018 Ivan Mahonin
6fa009
6fa009
    This program is free software: you can redistribute it and/or modify
6fa009
    it under the terms of the GNU General Public License as published by
6fa009
    the Free Software Foundation, either version 3 of the License, or
6fa009
    (at your option) any later version.
6fa009
6fa009
    This program is distributed in the hope that it will be useful,
6fa009
    but WITHOUT ANY WARRANTY; without even the implied warranty of
6fa009
    MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
6fa009
    GNU General Public License for more details.
6fa009
6fa009
    You should have received a copy of the GNU General Public License
6fa009
    along with this program.  If not, see <http: licenses="" www.gnu.org="">.</http:>
6fa009
*/
6fa009
6fa009
#include <cassert></cassert>
6fa009
6fa009
#include <algorithm></algorithm>
6fa009
#include <iostream></iostream>
6fa009
6fa009
#include "cudarender.h"
6fa009
#include "measure.h"
6fa009
6fa009
6fa009
using namespace std;
6fa009
6fa009
6fa009
CudaRender::CudaRender(CudaContext &cu):
6fa009
	cu(cu),
6fa009
	contour_module(),
6fa009
	contour_path_kernel(),
6fa009
	contour_fill_kernel(),
6fa009
	surface(),
6fa009
	points_buffer(),
6fa009
	mark_buffer(),
6fa009
	surface_image()
6fa009
{
6fa009
	cu.err = cuModuleLoad(&contour_module, "cuda/contour.ptx");
6fa009
	assert(!cu.err);
6fa009
a7622f
	//cu.err = cuModuleGetFunction(&contour_path_kernel, contour_module, "path");
a7622f
	cu.err = cuModuleGetFunction(&contour_path_kernel, contour_module, "_Z4pathiiPiPK6float2iii");
6fa009
	assert(!cu.err);
6fa009
a7622f
	//cu.err = cuModuleGetFunction(&contour_fill_kernel, contour_module, "fill");
a7622f
	cu.err = cuModuleGetFunction(&contour_fill_kernel, contour_module, "_Z4filliP4int2P6float4S1_4int4");
6fa009
	assert(!cu.err);
6fa009
}
6fa009
6fa009
CudaRender::~CudaRender() {
6fa009
	send_points(NULL, 0);
6fa009
	send_surface(NULL);
6fa009
6fa009
	cu.err = cuModuleUnload(contour_module);
6fa009
	assert(!cu.err);
6fa009
}
6fa009
6fa009
void CudaRender::send_surface(Surface *surface) {
6fa009
	if (this->surface) {
6fa009
		wait();
6fa009
6fa009
		cu.err = cuMemFree(surface_image);
6fa009
		assert(!cu.err);
6fa009
		surface_image = 0;
6fa009
6fa009
		cu.err = cuMemFree(mark_buffer);
6fa009
		assert(!cu.err);
6fa009
		mark_buffer = 0;
6fa009
	}
6fa009
6fa009
	this->surface = surface;
6fa009
6fa009
	if (this->surface) {
6fa009
		cu.err = cuMemAlloc(&surface_image, surface->data_size());
6fa009
		assert(!cu.err);
6fa009
6fa009
		cu.err = cuMemcpyHtoD(surface_image, surface->data, surface->data_size());
6fa009
		assert(!cu.err);
6fa009
a7622f
		cu.err = cuMemAlloc(&mark_buffer, surface->count()*sizeof(vec2i));
6fa009
		assert(!cu.err);
6fa009
a7622f
		cu.err = cuMemsetD32(mark_buffer, 0, 2*surface->count());
6fa009
		assert(!cu.err);
6fa009
	}
6fa009
}
6fa009
6fa009
Surface* CudaRender::receive_surface() {
6fa009
	if (surface) {
6fa009
		wait();
6fa009
		cu.err = cuMemcpyDtoH(surface->data, surface_image, surface->data_size());
6fa009
		assert(!cu.err);
6fa009
	}
6fa009
	return surface;
6fa009
}
6fa009
6fa009
void CudaRender::send_points(const vec2f *points, int count) {
6fa009
	if (points_buffer) {
6fa009
		wait();
6fa009
		cu.err = cuMemFree(points_buffer);
6fa009
		assert(!cu.err);
6fa009
		points_buffer = 0;
6fa009
	}
6fa009
6fa009
	if (points && count > 0) {
6fa009
		cu.err = cuMemAlloc(&points_buffer, count*sizeof(vec2f));
6fa009
		assert(!cu.err);
6fa009
6fa009
		cu.err = cuMemcpyHtoD(points_buffer, points, count*sizeof(vec2f));
6fa009
		assert(!cu.err);
6fa009
	}
6fa009
}
6fa009
6fa009
void CudaRender::draw(const Path &path) {
6fa009
	assert(surface);
6fa009
	assert(points_buffer);
6fa009
6fa009
	ContextRect bounds;
6fa009
	bounds.minx = max(1, path.bounds.minx);
6fa009
	bounds.maxx = min(surface->width, path.bounds.maxx);
6fa009
	bounds.miny = max(0, path.bounds.miny);
6fa009
	bounds.maxy = min(surface->height, path.bounds.maxy);
6fa009
	if ( bounds.minx >= bounds.maxx
6fa009
	  || bounds.miny >= bounds.maxy
6fa009
	  || path.begin >= path.end ) return;
6fa009
6fa009
	vec2i boundsx(bounds.minx, bounds.maxx);
6fa009
6fa009
	size_t group_size, count;
6fa009
6fa009
	count = path.end - path.begin;
a7622f
	group_size = 128;
6fa009
6fa009
	count = (count - 1)/group_size + 1;
6fa009
	cu.err = cuLaunchKernel(
6fa009
		contour_path_kernel,
6fa009
		count, 1, 1,
6fa009
		group_size, 1, 1,
6fa009
		0, 0, 0,
6fa009
		CudaParams()
6fa009
			.add(surface->width)
6fa009
			.add(surface->height)
6fa009
			.add(mark_buffer)
6fa009
			.add(points_buffer)
6fa009
			.add(path.begin)
6fa009
			.add(path.end)
6fa009
			.add(bounds.minx)
6fa009
			.get_extra() );
6fa009
	assert(!cu.err);
6fa009
a7622f
	count = bounds.maxx - bounds.minx;
a7622f
	group_size = 16;
6fa009
6fa009
	count = (count - 1)/group_size + 1;
6fa009
	cu.err = cuLaunchKernel(
6fa009
		contour_fill_kernel,
6fa009
		count, 1, 1,
6fa009
		group_size, 1, 1,
6fa009
		0, 0, 0,
6fa009
		CudaParams()
6fa009
			.add(surface->width)
6fa009
			.add(mark_buffer)
6fa009
			.add(surface_image)
6fa009
			.add(path.color, 16)
6fa009
			.add(bounds, 16)
6fa009
			.get_extra() );
6fa009
	assert(!cu.err);
6fa009
}
6fa009
6fa009
void CudaRender::wait() {
6fa009
	cu.err = cuStreamSynchronize(0);
6fa009
	assert(!cu.err);
6fa009
}
6fa009