binarize/binarize.cu

   1 /*
   2         Copyright (C) 2012  fmaj7b5.info
   3
   4         This program is free software: you can redistribute it and/or modify
   5         it under the terms of the GNU General Public License as published by
   6         the Free Software Foundation, either version 2 of the License, or
   7         (at your option) any later version.
   8
   9         This program is distributed in the hope that it will be useful,
  10         but WITHOUT ANY WARRANTY; without even the implied warranty of
  11         MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
  12         GNU General Public License for more details.
  13
  14         You should have received a copy of the GNU General Public License
  15         along with this program.  If not, see <http://www.gnu.org/licenses/>.
  16 */
  17
  18 #include "stdafx.h"
  19 #include <cuda.h>
  20
  21 #include "binarize.cuh"
  22 #include "DeviceMemory.cuh"
  23
  24 using namespace FM7b5;
  25
  26 void
  27 FM7b5::binarize_gpu(ImageGray& out, const ImageGray& in, const uint8_t thres)
  28 {
  29         if (in.width() != out.width() || in.height() != out.height()) {
  30                 throw std::runtime_error("sizes of input and output images are diferent.");
  31         }
  32
  33         const size_t width(in.width()), height(in.height()), bpp(in.bpp());
  34
  35         const size_t threads_per_dim(32);
  36         dim3 threads_per_block(threads_per_dim, threads_per_dim);
  37         dim3 blocks_per_grid((width + threads_per_block.x - 1)/ threads_per_block.x,
  38                              (height + threads_per_block.y - 1)/ threads_per_block.y);
  39
  40         // allocate input/output memories
  41         memory::LinearPitch<uint8_t> d_in(width, height), d_out(width, height);
  42
  43         // copy an input image to device memory
  44         d_in.copy_from(in.data(), bpp * width, height, in.stride());
  45
  46         // launch kernel
  47         binarize<<<blocks_per_grid, threads_per_block>>>(d_out.ref(), d_in.ref(), width, height, thres);
  48
  49         // copy the result back to host memory
  50         d_out.copy_to(out.data(), bpp * width, height, out.stride());
  51 }