NPP : Canny変換

NPP(NVIDIA Performance Primitive) の中に Canny変換 を見つけました。
どうやら CUDA 8.0 で新たに追加されたみたいです。

Canny変換は画像の輪郭を抽出するもので、Sobel/Scharr変換よりシャープな輪郭線を描いてくれます。 Sobel/Scharr変換で得られた輝度勾配の稜線を見つけてくれるってゆーか。

早速試してみました。 OpenCV 3.1 を使ってWeb-cameraからの画像のキャプチャと描画を行います。ダンドリはこんな。

1. Web-cameraからキャプチャした画像を
2. Device-memoryにコピー
3. モノクロ化し
4. Canny変換を施します。
5. 変換後のモノクロ画像をカラー化(RGBに同じ値を入れるだけ)し
6. Hostに書き戻して
7. 描画!

/*
 * DO NOT FORGET nvcc option : --expt-extended-lambda
 */

// std
#include <iostream>

// OpenCV
#include <opencv2/opencv.hpp>

// CUDA
#include <cuda_runtime.h>
#include <device_launch_parameters.h>
#include <npp.h>

// カーネル関数 二次元のtransform 
//    dst[y][x] = fun(src[y][x]) 
//       where : 0 <= x < width, 0 <= y < height
template<typename T, typename U, typename Function>
__global__ void kernel_transform2D(unsigned int  width, unsigned int height, 
                                        const T* src,         size_t src_pitch,
                                              U* dst,         size_t dst_pitch,
                                       Function  fun) {
  unsigned int x = blockDim.x * blockIdx.x + threadIdx.x;
  unsigned int y = blockDim.y * blockIdx.y + threadIdx.y;
  if ( x < width && y < height ) {
    ((U*)((char*)dst + dst_pitch*y))[x] = fun(((const T*)((const char*)src + src_pitch*y))[x]);
  }
}

void color2gray(unsigned int  width, unsigned int height, 
                      uchar3* src,         size_t src_pitch,
                       uchar* dst,         size_t dst_pitch) {
  kernel_transform2D<<<dim3((width+31)/32, (height+7)/8), dim3(32,8)>>>(
    width, height, 
    src, src_pitch, 
    dst, dst_pitch,
    [] __device__ (const uchar3 v) -> uchar { 
       int t = (v.x + v.y*7 + v.z*2)/10; 
       if ( t <   0 ) t = 0; 
       if ( t > 255 ) t = 255; 
       return (uchar)t; 
    }
  );
}

void gray2color(unsigned int  width, unsigned int height, 
                       uchar* src,         size_t src_pitch,
                      uchar3* dst,         size_t dst_pitch) {
  kernel_transform2D<<<dim3((width+31)/32, (height+7)/8), dim3(32,8)>>>(
    width, height, 
    src, src_pitch, 
    dst, dst_pitch,
    [] __device__ (const uchar v) -> uchar3 { return make_uchar3(v,v,v); }
  );
}

int main(int argc, char *argv[]) {
  cv::VideoCapture camera(0);

  cv::namedWindow("original", CV_WINDOW_AUTOSIZE);
  cv::namedWindow("canny", CV_WINDOW_AUTOSIZE);

  cv::Mat frame;
  cv::Mat canny;

  uchar3* d_frame;
  uchar*  d_gray;
  uchar*  d_canny;
  size_t  d_frame_pitch;
  size_t  d_gray_pitch;
  size_t  d_canny_pitch;
  Npp8u*  d_buffer;
  NppiSize size;

 
  // 一発目のキャプチャでフレームのサイズがわかるから
  // (そして多分その後ずっと変わらんだろから)
  // それを基にdevice-memoryを確保
  camera >> frame;
  size.width = (int)frame.size().width;
  size.height = (int)frame.size().height;

  cudaMallocPitch(&d_frame, &d_frame_pitch, size.width*sizeof(uchar3), size.height);
  cudaMallocPitch(&d_gray,  &d_gray_pitch,  size.width,                size.height);
  cudaMallocPitch(&d_canny, &d_canny_pitch, size.width,                size.height);

  // Cannyに引き渡すパラメータ
  NppiSize  nroi    = size;
  NppiPoint noffset = { 0, 0 };
  Npp16s    nlow_threshold  = 50;  // これと
  Npp16s    nhigh_threshold = 150; // これは適宜調整。

  // Canny変換に必要なバッファを確保
  {
  int buffer_size;
  nppiFilterCannyBorderGetBufferSize(size, &buffer_size);
  cudaMalloc(&d_buffer, buffer_size);
  }

  canny = frame.clone();
  std::cout 
    << "width,height   = " << size.width << ',' << size.height  
    << "\nstep           = " << frame.step 
    << "\ndepth, channel = " << frame.depth() << ',' << frame.channels()
    << "\n***** [ESC] to exit. *****\n";

  while ( cv::waitKey(10) != 0x1b ) {

    // [1] 画像を frame にキャプチャ
    camera >> frame;
    cv::imshow("original", frame);

    // [2] frame から d_frame へコピー
    cudaMemcpy2D(d_frame, d_frame_pitch, frame.data, frame.step, 
                 size.width*sizeof(uchar3), size.height, cudaMemcpyDefault);
    // [3] d_frame をモノクロ化して d_gray へ
    color2gray(size.width, size.height, d_frame, d_frame_pitch, d_gray, d_gray_pitch);

    // [4] d_gray に Canny変換カマして d_canny へ
    nppiFilterCannyBorder_8u_C1R(d_gray,  (int)d_gray_pitch,  size, noffset,
                                 d_canny, (int)d_canny_pitch, nroi,
                                 NPP_FILTER_SOBEL, NPP_MASK_SIZE_3_X_3,
                                 nlow_threshold, nhigh_threshold,
                                 nppiNormL2, NPP_BORDER_REPLICATE, 
                                 d_buffer);

    // [5] d_canny をカラー化(RGBを同じ値にするだけ)して d_frame へ
    gray2color(size.width, size.height, d_canny, d_canny_pitch, d_frame, d_frame_pitch);

    // [6] d_frame を canny へコピー
    cudaMemcpy2D(canny.data, canny.step, d_frame, d_frame_pitch, 
                 size.width*sizeof(uchar3), size.height, cudaMemcpyDefault);

    // [7] 描画!
    cv::imshow("canny", canny);
  }

  // あとしまつ
  cudaFree(d_frame);
  cudaFree(d_gray);
  cudaFree(d_canny);
  cudaFree(d_buffer);
}

こんな輪郭線を描いてくれます;