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;
}
|
執行結果:
留言列表
雜項記錄 (3)
