OpenCV和Cuda结合编程

http://www.cnblogs.com/dwdxdy/p/3528711.html

一、利用OpenCV中提供的GPU模块

缺点是受限于OpenCV库的发展和更新，当需要完成一些自定义的操作时（OpenCV中没有提供相应的库），难以满足应用的需求，需要自己实现自定义操作的并行实现。此外，针对一些特殊需求，OpenCV提供并行处理函数，其性能优化并不是最优的，在具体的应用时，可能需要进一步优化，提高性能。

二、单独使用Cuda API编程

缺点是使用复杂，代码编写量较多，需要熟悉Cuda相关资料和API接口
三、利用OpenCV中提供接口，并结合Cuda API编程

利用OpenCV已经提供的部分接口，完成一些Cuda编程的基本处理，简化编程的复杂程度；只是根据自己业务需求，自定义内核函数或扩展OpenCV已提供的内核函数。这样既可以充分利用OpenCV的特性，又可以满足业务的不同需求，使用方便，且易于扩展。

 

//main.cpp
#include <iostream>
#include <opencv2/opencv.hpp>
#include <opencv2/gpu/gpu.hpp>
using namespace cv;
using namespace cv::gpu;
void swap_rb(const GpuMat& src,GpuMat& dst,Stream& stream = Stream::Null());
int main()
{
Mat image = imread("lena.jpg");
imshow("src",image);
GpuMat gpuMat,output;

gpuMat.upload(image);
swap_rb(gpuMat,output);
output.download(image);

imshow("gpu",image);
waitKey(0);
return 0;
}

 

//swap_rb.cpp     定义了并行操作的入口函数，
#include <opencv2/gpu/gpu.hpp>
#include <opencv2/gpu/stream_accessor.hpp>
using namespace cv;
using namespace cv::gpu;

void swap_rb_caller(const PtrStepSz<uchar3>& src,PtrStep<uchar3> dst,cudaStream_t stream);

void swap_rb(const GpuMat& src,GpuMat& dst,Stream& stream = Stream::Null())//供主函数调用
{
    CV_Assert(src.type() == CV_8UC3);
    dst.create(src.size(),src.type());
    cudaStream_t s = StreamAccessor::getStream(stream);
    swap_rb_caller(src,dst,s);//GpuMat    -》   PtrStepSz<uchar3>

}

 

//swap_rb.cu
#include <opencv2/core/cuda_devptrs.hpp>
using namespace cv;
using namespace cv::gpu;
//自定义内核函数
__global__ void swap_rb_kernel(const PtrStepSz<uchar3> src,PtrStep<uchar3> dst)
{
int x = threadIdx.x + blockIdx.x * blockDim.x;
int y = threadIdx.y + blockIdx.y * blockDim.y;

if(x < src.cols && y < src.rows)
{
uchar3 v = src(y,x);
dst(y,x) = make_uchar3(v.z,v.y,v.x);
}
}

 

 

-----------------------------------------------------------------------------------------------------------------

全程自己编写

__global__ void swap_rb_kernel(const uchar3* src,uchar3* dst,int width,int height)
{
    int x = threadIdx.x + blockIdx.x * blockDim.x;
    int y = threadIdx.x + blockIdx.y * blockDim.y;
    
    if(x < width && y < height)
    {
        uchar3 v = src[y * width + x];
        dst[y * width + x].x = v.z;//PtrStep<uchar3>       交换(x, y)     变成(y, x)
        dst[y * width + x].y = v.y;
        dst[y * width + x].z = v.x;
    }
}
若image.data包含字节对齐的空白数据，上述程序无法完成正常的处理操作
void swap_rb_caller(const uchar3* src,uchar3* dst,int width,int height)
{
    dim3 block(32,8);
    dim3 grid((width + block.x - 1)/block.x,(height + block.y - 1)/block.y);
    
    swap_rb_kernel<<<grid,block,0>>>(src,dst,width,height);
    cudaThreadSynchronize();
}

int main()
{
    Mat image = imread("lena.jpg");
    imshow("src",image);
    
    size_t memSize = image.cols*image.rows*sizeof(uchar3);
    uchar3* d_src = NULL;
    uchar3* d_dst = NULL;
    CUDA_SAFE_CALL(cudaMalloc((void**)&d_src,memSize));
    CUDA_SAFE_CALL(cudaMalloc((void**)&d_dst,memSize));
    CUDA_SAFE_CALL(cudaMempcy(d_src,image.data,memSize,cudaMemcpyHostToDevice));
    
    swap_rb_caller(d_src,d_dst,image.cols,image.rows);
    
    CUDA_SAFE_CALL(cudaMempcy(image.data,d_dst,memSize,cudaMemcpyDeviceToHost));
    imshow("gpu",image);
    waitKey(0);
    
    CUDA_SAFE_CALL(cudaFree(d_src));
    CUDA_SAFE_CALL(cudaFree(d_dst));
    return 0;
}

void swap_rb_caller(const PtrStepSz<uchar3>& src,PtrStep<uchar3> dst,cudaStream_t stream)//被swap_rb调用

{

dim3 block(32,8);//?
dim3 grid((src.cols + block.x - 1)/block.x,(src.rows + block.y - 1)/block.y);

swap_rb_kernel<<<grid,block,0,stream>>>(src,dst);// PtrStep<uchar3>  ->dim3
if(stream == 0)
cudaDeviceSynchronize();
}

 

　　swap_rb.cu文件定义了内核函数和内核函数的调用函数，在调用函数中，设置内核函数的调用参数。

　　swap_rb.cpp文件定义了并行操作的入口函数，即主程序完成并行操作的需要调用的函数，其主要是封装内核函数的调用函数，并添加输入参数的验证、根据输入参数选择不同内核函数等操作。

　　main.cpp文件主程序，完成数据的输入、业务的处理和数据的输出。