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

　　目前，OpenCV中已提供了许多GPU函数，直接使用OpenCV提供的GPU模块，可以完成大部分图像处理的加速操作。

　　基本使用方法，请参考：

　　该方法的优点是使用简单，利用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<<
     
      >>(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 
     
      using namespace cv;using namespace cv::gpu;//自定义内核函数__global__ void swap_rb_kernel(const PtrStepSz
      
        src,PtrStep
       
         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
        
         & src,PtrStep
         
           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<<
          
           >>(src,dst); if(stream == 0) cudaDeviceSynchronize();}
          
         
        
       
      
     

//swap_rb.cpp#include 
     
      #include 
      
       using namespace cv;using namespace cv::gpu;void swap_rb_caller(const PtrStepSz
       
        & src,PtrStep
        
          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 
     
      #include 
      
       #include 
       
        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文件主程序，完成数据的输入、业务的处理和数据的输出。

总结

　　编程简易性和可控性是相对的，编程越方便，就越不容易控制。实际应用过程中，应当寻求编程简易性和可控性的平衡点，应根据应用需求，选取适当的方法，一般建议采用方法三。