【OpenCV & CUDA】OpenCV和Cuda結合編程

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

　　目前，OpenCV中已提供了許多GPU函數，直接使用OpenCV提供的GPU模塊，可以完成大部分圖像處理的加速操作。

　　基本使用方法，請參考：http://www.cnblogs.com/dwdxdy/p/3244508.html

　　該方法的優點是使用簡單，利用GpuMat管理CPU與GPU之間的數據傳輸，而且不需要關注內核函數調用參數的設置，使用過程中，只需要關注處理的邏輯操作。

　　缺點是受限於OpenCV庫的發展和更新，當需要完成一些自定義的操作時（OpenCV中沒有提供相應的庫），難以滿足應用的需求，需要自己實現自定義操作的並行實現。此外，針對一些特殊需求，OpenCV提供並行處理函數，其性能優化並不是最優的，在具體的應用時，可能需要進一步優化，提高性能。

二、單獨使用Cuda API編程

　　利用Cuda Runtime API、Cuda Driver API實現一些操作的並行加速，使用過程需要管理CPU與GPU之間的數據傳輸，內核函數調用參數的設置，內核函數的優化等。

　　優點是處理過程受控於用戶，用戶可以實現更多的並行加速處理操作。

　　缺點是使用復雜，代碼編寫量較多，需要熟悉Cuda相關資料和API接口。下面是簡單的示例程序：

__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;
        dst[y * width + x].y = v.y;
        dst[y * width + x].z = v.x;
    }
}

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;
}

　　上述代碼中，使用cudaMalloc，cudaMemcpy，cudaFree管理內存的分配、傳輸和釋放。

　　注意：若image.data包含字節對齊的空白數據，上述程序無法完成正常的處理操作。

三、利用OpenCV中提供接口，並結合Cuda API編程

　　利用OpenCV已經提供的部分接口，完成一些Cuda編程的基本處理，簡化編程的復雜程度；只是根據自己業務需求，自定義內核函數或擴展OpenCV已提供的內核函數。這樣既可以充分利用OpenCV的特性，又可以滿足業務的不同需求，使用方便，且易於擴展。下面是簡單的示例程序：

//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);
    }
}

void swap_rb_caller(const PtrStepSz<uchar3>& src,PtrStep<uchar3> dst,cudaStream_t stream)
{
    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);
    if(stream == 0)
        cudaDeviceSynchronize();
}

//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);
}

//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.cu文件定義了內核函數和內核函數的調用函數，在調用函數中，設置內核函數的調用參數。

　　swap_rb.cpp文件定義了並行操作的入口函數，即主程序完成並行操作的需要調用的函數，其主要是封裝內核函數的調用函數，並添加輸入參數的驗證、根據輸入參數選擇不同內核函數等操作。

　　main.cpp文件主程序，完成數據的輸入、業務的處理和數據的輸出。

總結

　　編程簡易性和可控性是相對的，編程越方便，就越不容易控制。實際應用過程中，應當尋求編程簡易性和可控性的平衡點，應根據應用需求，選取適當的方法，一般建議采用方法三。