/*
	Copyright (C) 2012  fmaj7b5.info

	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 2 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 "stdafx.h"
#include <cuda.h>

#include "binarize.cuh"
#include "DeviceMemory.cuh"

using namespace FM7b5;

void
FM7b5::binarize_gpu(ImageGray& out, const ImageGray& in, const uint_fast8_t thres)
{
	if (in.width() != out.width() || in.height() != out.height()) {
		throw std::runtime_error("sizes of input and output images are diferent.");
	}

	const size_t width(in.width()), height(in.height()), bpp(in.bpp());

	const size_t threads_per_dim(32);
	dim3 threads_per_block(threads_per_dim, threads_per_dim);
	dim3 blocks_per_grid((width + threads_per_block.x - 1)/ threads_per_block.x,
	                     (height + threads_per_block.y - 1)/ threads_per_block.y);

	// allocate input/output memories
	memory::LinearPitch<uint_fast8_t> d_in(width, height), d_out(width, height);

	// copy an input image to device memory
	d_in.copy_from(in.data(), bpp * width, height, in.stride());

	// launch kernel
	binarize<<<blocks_per_grid, threads_per_block>>>(d_out.ref(), d_in.ref(), width, height, thres);

	// copy the result back to host memory
	d_out.copy_to(out.data(), bpp * width, height, out.stride());
}