[CUDA+OpenCV] GpuMat 結合 kernel 運算

OpenCV 的 GpuMat 類別, 可以在 device 的 global memory 建立 Mat 物件, 可以被 kernel 函數直接存取, 以下以 sobel 運算子為例, 將 srcMat (位於host memory) 載入影像後, 轉換至 srcGMat (位於device memory), 經處理後放至 destGMat (位於device memory), 然後下載回 destMat (位於host memory)再顯示出來.
程式中使用 OpenCV 定義於 cuda_types.hpp 的 PtrStep 的樣板結構, 代入定義於 WIndows.h 的 byte, 這樣比較方便. 如果要用 cuda 編譯器 nvcc 內建的  uchar1 , 查看 vector_types.h, 它是一個結構 (向量):

struct __device_builtin__ uchar1 {     unsigned char x; };

這樣無法直接使用 + = * / 等之數學運算子, 亦無法進行 type cast, 必須使用 .x 成員才能存取, 例如 :

__global__ void sobel2Kernel(PtrStep<uchar1> dest, PtrStep<uchar1> src, int width, int height) {      int x = threadIdx.x; // +blockIdx.x * blockDim.x;      int y = blockIdx.x; // threadIdx.y + blockIdx.y * blockDim.y;      if (x > 0 && x < width - 1 && y>0 && y < height - 1)      {          int Gx = (int)src(y - 1, x - 1).x - src(y - 1, x + 1).x               + 2 * src(y, x - 1).x - 2 * src(y, x + 1).x               + src(y + 1, x - 1).x - src(y + 1, x + 1).x; …………

這樣程式寫起來不方便, 因此本例中使用 byte 類別. 以下為完整 kernel.cu 程式碼 :

#include "cuda_runtime.h" #include "device_launch_parameters.h" #include <opencv2/core.hpp> #include <opencv2/imgcodecs.hpp> #include <opencv2/highgui.hpp> #include "opencv2/imgproc.hpp" #include <opencv2/core/cuda.hpp> #include <stdio.h> #include <Windows.h> using namespace cv; using namespace cuda; __global__ void sobelKernel(PtrStep<byte> dest, PtrStep<byte> src, int width, int height) {      int x = threadIdx.x; // +blockIdx.x * blockDim.x;      int y = blockIdx.x; // threadIdx.y + blockIdx.y * blockDim.y;      if (x>0 && x < width-1 && y>0 && y < height-1)      {          int Gx   = (int)src(y-1, x - 1) - src(y - 1, x + 1)                   + 2*src(y, x-1) - 2 * src(y, x+1)                   + src(y + 1, x - 1) - src(y + 1, x + 1);          int Gy = (int)src(y - 1, x - 1) - src(y + 1, x - 1)                   + 2 * src(y-1, x) - 2 * src(y+1, x)                   + src(y -1, x + 1) - src(y + 1, x + 1);          dest(y, x) = byte(sqrtf(Gx*Gx + Gy * Gy)); //sqrtf() is a device function      }      else      {          dest(x, y) = 0;      } } int main() {      char *filename = "D:\\LENA.JPG";      // load source image      Mat srcMat = imread(filename, IMREAD_GRAYSCALE); // Read the file and convert to grayscale      GpuMat  srcGMat(srcMat);      // allocate dest image buffer      GpuMat destGMat(srcGMat.size(), srcGMat.type());       //run processing kernel      int width = srcGMat.cols;      int height = srcGMat.rows;      sobelKernel << <height, width >> > (destGMat, srcGMat, width, height);      // get processing result      Mat destMat;      destGMat.download(destMat);      //show results      namedWindow("Source window", WINDOW_AUTOSIZE); // Create a window for display.      imshow("Source window", srcMat); // Show our image inside it.      namedWindow("Dest window", WINDOW_AUTOSIZE); // Create a window for display.      imshow("Dest window", destMat); // Show our image inside it.      waitKey(0); // Wait for a keystroke in the window      return 0; }

執行結果: