|
|
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 |
|