cuda与openCV结合编程（一）

    学习计算机图像处理算法的童鞋，就不得不学习cuda，为啥呢？因为图像处理一般都是矩阵运算，动不动就是百万的计算量这个时候优化计算时间是必不可少的。openCV本身提供了很多cuda函数，能够满足大多数用户的需求。但是也不绝对，有时候我们需要自己定义一个内核函数进行优化，当然你也可以用openGL或者多线程，openCV也提供较好的支持，掌握一种或多种加速算法，对程序员特别是算法工程师来讲很重要。
    闲话不多说，再学习了cuda的基础以后【cuda并行编程基础（一）、cuda并行编程基础（二）】，我们其实就具备与opencv联合编程的能力，虽然不是最佳优化，但已经可以满足大多数需求了。
一、cuda与openCV结合方法
    （下面仅涉及windows环境）
    1.我们知道，cuda代码一般以.cu结尾（windows，其他系统除外，下同），它的编译器是nvcc，编译时它会将CPU代码和GPU代码分开，CPU部分其实与gcc编译差不多，GPU部分就按照nvcc的规则编译，这玩意其实并不复杂;
    2.openCV的代码一般都是以.cpp结尾，它的编译器一般是gcc、g++(或者其他相似编译器),那么可不可以将openCV代码用nvcc编译呢?答案是肯定的,但在windows系统,你得把它改为.cu结尾.
    3.所以，在windows系统里面，你有两种办法让openCV结合cuda编程：
        a.openCV正常编译，cuda代码编译好后，作为静态库引入openCV调用；
        b.openCV与cuda代码混在一起，统一用nvcc编译。
二、如何编写代码
    (下面openCV基于3.2.0版本)
    openCV是一个非常强大的视觉算法库，当然也支持cuda咯。
    cv::cuda是一个专门处理cuda的命名空间，你在这个命名空间里面可以看到很多已经集成好的函数。
        如：cuda::remap()、cuda::add()等
    我们要用到的是cuda::PtrStepSz<T>的模板，以及cuda::GpuMat
    比如：如果我们有一个cuda::GpuMat类型的img，我们怎么传入cuda里面呢？答案就是，直接将img传到cuda::PtrStepSz里面，他们是不是等同，但是可以互传数据，具体见样例。至于传到cuda::PtrStepSz里面如何操作，那就跟cuda差不多了。
    除了cuda::PtrStepSz，openCV还有其他接口可以提供互传，自己去摸索啦，这里就不啰嗦了。

    至于cuda与openCV的结合编程效率问题？哈，谁用谁知道，你不用也无需知道，有兴趣自己去测一下咯，反正笔者是墙裂推荐的，后面有空再讲效率问题。
三、常见错误
    1.cudaErrorMemoryAllocation，主要是申请空间太大，超出了GPU限制；
    2.cudaErrorLaunchFailure，访问了非法地址，比如index超过了数组大小；
    3.cuda与vs2015结合编程，偶尔会出现抽筋的问题，比如你这次编译出错，改正了以后再编译还出错，建议要重新编译时，把以前的编译生成的东西全删掉，这样就保险多了，笔者遇见多次这种情况；
    4.<<<>>>内核符号报错，要确定它出现在cu文件里而不是cpp文件里，cu文件会显示红色，不用管它；
    5.静态库的编写规范，额，自己上网研究吧，其实我写得也不太规范，吐槽一下，网上好多技术文章抄来抄去很没意思，很多大牛又写得太过高深，研究不出个所以然来，也希望能够在各个层次都有合适的文章介绍吧，这样入门和进阶也不会太困难。

第一个程序，直接在cu文件实现cuda与opencv结合编程，非常重要哦
一般我们不这么使用，因为cuda作为独立的编程方式，放在一起容易混乱，而且为支持高速运算，一般都使用c运算，而不是c++

//opencv_cuda.cu:使用自定义函数，实现cuda版本图片翻转//authored by alpc40//version：visual studio 2015\cuda toolkit 8.0\openCV 3.2.0#include "cuda_runtime.h"#include "device_launch_parameters.h"#include <stdio.h>#include<opencv2/opencv.hpp>#include<iostream>using namespace std;using namespace cv;#ifdef _DEBUG#pragma comment ( lib,"opencv_core320d.lib")#pragma comment ( lib,"opencv_highgui320d.lib")#pragma comment ( lib,"opencv_calib3d320d.lib")#pragma comment ( lib,"opencv_imgcodecs320d.lib")#pragma comment ( lib,"opencv_imgproc320d.lib")#pragma comment ( lib,"opencv_cudaimgproc320d.lib")#pragma comment ( lib,"opencv_cudaarithm320d.lib")#pragma comment ( lib,"cudart.lib")#else#pragma comment ( lib,"opencv_core320.lib")#pragma comment ( lib,"opencv_highgui320.lib")#pragma comment ( lib,"opencv_calib3d320.lib")#pragma comment ( lib,"opencv_imgcodecs320.lib")#pragma comment ( lib,"opencv_imgproc320.lib")#pragma comment ( lib,"opencv_cudaimgproc320.lib")#pragma comment ( lib,"opencv_cudaarithm320.lib")#pragma comment ( lib,"cudart.lib")#endif//出错处理函数#define CHECK_ERROR(call){\    const cudaError_t err = call;\    if (err != cudaSuccess)\    {\        printf("Error:%s,%d,",__FILE__,__LINE__);\        printf("code:%d,reason:%s\n",err,cudaGetErrorString(err));\        exit(1);\    }\}//内核函数：实现上下翻转__global__ void swap_image_kernel(cuda::PtrStepSz<uchar3> cu_src, cuda::PtrStepSz<uchar3> cu_dst, int h, int w){    //计算的方法：参看前面两文    unsigned int x = blockDim.x * blockIdx.x + threadIdx.x;    unsigned int y = blockDim.y * blockIdx.y + threadIdx.y;    //为啥要这样限制：参看前面两文    if (x < cu_src.cols && y < cu_src.rows)    {        //为何不是h-y-1,而不是h-y，自己思考哦        cu_dst(y, x) = cu_src(h - y - 1, x);    }}//调用函数，主要处理block和grid的关系void swap_image(cuda::GpuMat src,cuda::GpuMat dst,int h, int w){    assert(src.cols == w && src.rows ==h);    int uint = 32;    //参考前面两文的block和grid的计算方法，注意不要超过GPU限制    dim3 block(uint, uint);    dim3 grid((w + block.x - 1) / block.x, (h + block.y - 1) / block.y);    printf("grid = %4d %4d %4d\n",grid.x,grid.y,grid.z);    printf("block= %4d %4d %4d\n",block.x,block.y,block.z);    swap_image_kernel << <grid, block >> > (src,dst,h,w);    //同步一下，因为计算量可能很大    CHECK_ERROR(cudaDeviceSynchronize());}int main(int argc,char **argv){    Mat src, dst;    cuda::GpuMat cu_src, cu_dst;    int h, w;    //根据argv[1]读入图片数据，BGR格式读进来    src = imread(argv[1]);    //检测是否正确读入    if (src.data == NULL)    {        cout << "Read image error" << endl;        return -1;    }    h = src.rows; w = src.cols;    cout <<"图片高：" << h << ",图片宽：" << w << endl;    //上传CPU图像数据到GPU，跟cudaMalloc和cudaMemcpy很像哦，其实upload里面就是这么写的    cu_src.upload(src);    //申请GPU空间，也可以到函数里申请，不管怎样总要申请，要不然内核函数会爆掉哦    cu_dst = cuda::GpuMat(h, w, CV_8UC3, Scalar(0, 0, 0));    //申请CPU空间    dst = Mat(h, w, CV_8UC3, Scalar(0, 0, 0));    //调用函数swap_image,由该函数调用内核函数，这样层次分明，不容易出错    //当然你也可以直接在这里调用内核函数，东西太多代码容易乱    swap_image(cu_src,cu_dst,h, w);    //下载GPU数据到CPU，与upload()对应    cu_dst.download(dst);    //显示cpu图像，如果安装了openCV集成了openGL,那可以直接显示GpuMat    imshow("dst",dst);    //等待按键    waitKey();    //写图片到文件    if(argc==3)        imwrite(argv[2],dst);    return 0;}

第二个程序，使用静态库的方式实现cuda与openCV的结合，非常重要哦
这种方式两相分离，更好实现了这种功能

//swap_image.cu:生成swap_image.lib，供主函数调用

//authored by alpc40//version：visual studio 2015\cuda toolkit 8.0\openCV 3.2.0#include "cuda_runtime.h"#include "device_launch_parameters.h"#include <stdio.h>#include<opencv2/opencv.hpp>using namespace cv;//出错处理函数#define CHECK_ERROR(call){\    const cudaError_t err = call;\    if (err != cudaSuccess)\    {\        printf("Error:%s,%d,",__FILE__,__LINE__);\        printf("code:%d,reason:%s\n",err,cudaGetErrorString(err));\        exit(1);\    }\}//内核函数：实现上下翻转__global__ void swap_image_kernel(cuda::PtrStepSz<uchar3> cu_src, cuda::PtrStepSz<uchar3> cu_dst, int h, int w){    //计算的方法：参看前面两文    unsigned int x = blockDim.x * blockIdx.x + threadIdx.x;    unsigned int y = blockDim.y * blockIdx.y + threadIdx.y;    //为啥要这样限制：参看前面两文    if (x < cu_src.cols && y < cu_src.rows)    {        //为何不是h-y-1,而不是h-y，自己思考哦        cu_dst(y, x) = cu_src(h - y - 1, x);    }}//调用函数，主要处理block和grid的关系，注意extern哦，它是库文件编写规范extern "C" void swap_image(cuda::GpuMat src, cuda::GpuMat dst, int h, int w){    assert(src.cols == w && src.rows == h);    int uint = 32;    //参考前面两文的block和grid的计算方法，注意不要超过GPU限制    dim3 block(uint, uint);    dim3 grid((w + block.x - 1) / block.x, (h + block.y - 1) / block.y);    printf("grid = %4d %4d %4d\n", grid.x, grid.y, grid.z);    printf("block= %4d %4d %4d\n", block.x, block.y, block.z);    swap_image_kernel << <grid, block >> > (src, dst, h, w);    //同步一下，因为计算量可能很大    CHECK_ERROR(cudaDeviceSynchronize());}

//opencv_cuda.cpp:第二个程序主函数,使用自编静态库，实现cuda版本图片翻转//authored by alpc40//version：visual studio 2015\cuda toolkit 8.0\openCV 3.2.0#include <stdio.h>#include<opencv2/opencv.hpp>#include<iostream>using namespace std;using namespace cv;#ifdef _DEBUG#pragma comment ( lib,"opencv_core320d.lib")#pragma comment ( lib,"opencv_highgui320d.lib")#pragma comment ( lib,"opencv_calib3d320d.lib")#pragma comment ( lib,"opencv_imgcodecs320d.lib")#pragma comment ( lib,"opencv_imgproc320d.lib")#pragma comment ( lib,"opencv_cudaimgproc320d.lib")#pragma comment ( lib,"opencv_cudaarithm320d.lib")#pragma comment ( lib,"cudart.lib")#pragma comment ( lib,"swap_image.lib")//别忘了加库#else#pragma comment ( lib,"opencv_core320.lib")#pragma comment ( lib,"opencv_highgui320.lib")#pragma comment ( lib,"opencv_calib3d320.lib")#pragma comment ( lib,"opencv_imgcodecs320.lib")#pragma comment ( lib,"opencv_imgproc320.lib")#pragma comment ( lib,"opencv_cudaimgproc320.lib")#pragma comment ( lib,"opencv_cudaarithm320.lib")#pragma comment ( lib,"cudart.lib")#pragma comment ( lib,"swap_image.lib")//别忘了加库#endif//这个声明很重要,调用静态库extern "C" void swap_image(cuda::GpuMat src,cuda::GpuMat dst,int w,int h);int main(int argc, char **argv){    Mat src, dst;    cuda::GpuMat cu_src, cu_dst;    int h, w;    //根据argv[1]读入图片数据，BGR格式读进来    src = imread(argv[1]);    //检测是否正确读入    if (src.data == NULL)    {        cout << "Read image error" << endl;        return -1;    }    h = src.rows; w = src.cols;    cout << "图片高：" << h << ",图片宽：" << w << endl;    //上传CPU图像数据到GPU，跟cudaMalloc和cudaMemcpy很像哦，其实upload里面就是这么写的    cu_src.upload(src);    //申请GPU空间，也可以到函数里申请，不管怎样总要申请，要不然内核函数会爆掉哦    cu_dst = cuda::GpuMat(h, w, CV_8UC3, Scalar(0, 0, 0));    //申请CPU空间    dst = Mat(h, w, CV_8UC3, Scalar(0, 0, 0));    //调用函数swap_image,由该函数调用内核函数，这样层次分明，不容易出错    //当然你也可以直接在这里调用内核函数，东西太多代码容易乱    swap_image(cu_src, cu_dst, h, w);    //下载GPU数据到CPU，与upload()对应    cu_dst.download(dst);    //显示cpu图像，如果安装了openCV集成了openGL,那可以直接显示GpuMat    imshow("dst", dst);    //等待按键    waitKey();    //写图片到文件    if (argc == 3)        imwrite(argv[2], dst);    return 0;}

PS：以下为图片效果，顺便为我家赣南脐橙打个广告