CUDA编程：与OpenCV结合

最新推荐文章于 2024-09-20 13:43:17 发布
最新推荐文章于 2024-09-20 13:43:17 发布
阅读量3.4k 收藏 14
点赞数 1
分类专栏： ------CUDA------ 文章标签： CUDA
版权声明：本文为博主原创文章，遵循 CC 4.0 BY-SA 版权协议，转载请附上原文出处链接和本声明。
本文链接：https://blog.csdn.net/Aidam_Bo/article/details/89176689
版权

前言

       学习计算机图像处理算法的童鞋，就不得不学习cuda，为啥呢？因为图像处理一般都是矩阵运算，动不动就是百万的计算量这个时候优化计算时间是必不可少的。openCV本身提供了很多cuda函数，能够满足大多数用户的需求。但是也不绝对，有时候我们需要自己定义一个内核函数进行优化，当然你也可以用openGL或者多线程，openCV也提供较好的支持，掌握一种或多种加速算法，对程序员特别是算法工程师来讲很重要。
       闲话不多说，在学习了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;
}


上图：

 对比下，两张图

ooMelloo
关注
专栏目录
使用Thrust库在OpenCV中编程实现cv::cuda::GpuMat
2301_79326559的博客
09-07 153
在计算机视觉和图像处理中，OpenCV是一个广泛使用的开源库，它提供了许多功能强大的图像处理和计算机视觉算法。OpenCV还提供了一个用于在GPU上执行图像处理操作的模块，即cv::cuda。为了更好地利用GPU的并行计算能力，我们可以使用Thrust库与OpenCV的cv::cuda::GpuMat结合使用。Thrust是一个用于并行计算的C++模板库，它提供了一组高级算法和数据结构，可简化在GPU上进行并行计算的过程。
利用CUDA与OpenCV实现高效图像处理：全面指南
qq_38334677的博客
06-05 458
CUDA（Compute Unified Device Architecture）是NVIDIA推出的一种并行计算架构，它使得GPU可以用于通用计算。相比于传统的CPU，GPU具有更多的处理核心，可以并行处理大量的数据。因此，在图像处理等需要大量计算的领域，利用CUDA能够显著提升处理速度。OpenCV（Open Source Computer Vision Library）是一个开源的计算机视觉库，包含了数百个图像处理和计算机视觉算法。
Python OpenCV with CUDA 开源项目教程
gitblog_00508的博客
08-23 243
Python OpenCV with CUDA 开源项目教程 python-opencv-cudacustom opencv_contrib module which exposes opencv cuda optical flow methods with python bindings项目地址:https://gitcode.com/gh_mirrors/py/python-opencv-c...
cuda与openCV结合编程（一）
热门推荐
alpc40计算机视觉开发
12-24 3万+
OpenCV与CUDA并行编程结合，是实现openCV优化时间的重要方法之一，搞图像处理的童鞋别错过哦
Opencv+Cuda编译的保姆级别教程
最新发布
m0_58326153的博客
09-20 1730
配置Opencv-cuda版本
opencv中使用cuda加速图像处理
xu.hyj
02-04 8801
opencv大多数只使用到了cpu的版本，实际上对于复杂的图像处理过程用cuda（特别是高分辨率的图像）可能会有加速效果。
OpenCV和Cuda结合编程
weixin_67271870的博客
12-09 1621
缺点是受限于OpenCV库的发展和更新，当需要完成一些自定义的操作时（OpenCV中没有提供相应的库），难以满足应用的需求，需要自己实现自定义操作的并行实现。swap_rb.cpp文件定义了并行操作的入口函数，即主程序完成并行操作的需要调用的函数，其主要是封装内核函数的调用函数，并添加输入参数的验证、根据输入参数选择不同内核函数等操作。该方法的优点是使用简单，利用GpuMat管理CPU与GPU之间的数据传输，而且不需要关注内核函数调用参数的设置，使用过程中，只需要关注处理的逻辑操作。只是根据自己业务需求，
opencv 使用cuda
AI.PLAY
03-31 3618
编译opencv gpu版 网上有各种使用cmake + vs编译opencv contirbute的教程，但编译总会出错，不过没关系，直接按照流程走一遍，大部分库都是会编译成功的，那些编译不成功的库一般也用不到。 随手附上一篇： https://blog.csdn.net/stjuliet/article/details/84640094 获得GPU信息 int getCu...
opencv+CUDA
xinxin_Tian的博客
03-22 471
1.https://blog.csdn.net/DumpDoctorWang/article/details/81032914 2.https://cloud.tencent.com/developer/article/1526401
【CUDA编程】opencv4 + CUDA 并行图像处理：图像均值滤波和图像反色
08-12
二、cuda与OpenCV结合方法 三、代码实例：图像均值滤波和图像反色 3.1 代码 3.2 代码说明 3.3 网格大小与线程块大小的确定 3.3.1 网格与线程块大小的限制 3.3.2 如何确定网格大小与线程块大小？ 3.4 并行与串行的...
opencv_cuda_opencvcuda_cuda+opencv_opencv_cudaopencv_cuda_源码.zip
10-05
1. **CUDA编程基础**：理解CUDA编程的基本概念，如线程块、网格、共享内存、全局内存等，是使用OpenCV CUDA模块的前提。开发者需要知道如何定义和管理CUDA内核，以及如何在CPU和GPU之间传输数据。 2. **OpenCV CUDA...
CUDA加速+VS+opencv
08-24
CUDA加速+VS+opencv
CUDA和OpenCV图像并行处理方法研究
04-29
CUDA和OpenCV图像并行处理方法研究，希望对研究并行计算的朋友们有帮助！
opencv与cuda混合编程模板
10-29
opencv与cuda混合编程，采用cuda里面的gpu模板与cuda结合，环境vs2008，opencv2.4.9.注意需要用Cmake重新编译opencv库
CUDA总结1——OpenCV与CUDA简介
windxgz的博客
09-11 5051
因为算法的需要，正常的CPU算法速度不够需要进行加速，OpenCV中正好加入了GPU计算的模块，OpenCV中有两种GPU的加速方式，一种是通用标准的opencl，另一种是NVIDIA的cuda加速。opencl是苹果公司提出的一种通用标准，多种平台支持的标准。cuda是NVIDIA提出的并行计算平台，只有NVIDIA的显卡支持，但是比起opencl，cuda在使用上更方便，因为opencl是开放标准，封装上做的可能没有那么好，现在NVIDIA的显卡使用的多一些，所以cuda也就比较多用。 使用OpenCV
Opencv + cuda 混编模式
slamer的专栏
10-24 1529
利用OpenCV中提供接口，并结合Cuda API编程 　　利用OpenCV已经提供的部分接口，完成一些Cuda编程的基本处理，简化编程的复杂程度；只是根据自己业务需求，自定义内核函数或扩展OpenCV已提供的内核函数。这样既可以充分利用OpenCV的特性，又可以满足业务的不同需求，使用方便，且易于扩展。下面是简单的示例程序： //swap_rb.cu #include using
OpenCV中的CUDA模块编程详解
CodeLancerX的博客
09-14 456
在计算机视觉和图像处理领域，OpenCV是一个非常流行的开源库，它提供了丰富的函数和算法来处理图像和视频数据。为了加速图像处理任务，OpenCV还提供了一个CUDA模块，该模块利用了NVIDIA的CUDA平台和GPU（图形处理器）进行并行计算。本文将详细介绍OpenCV中的CUDA模块编程，并提供相应的源代码示例。通过使用CUDA模块，我们可以利用GPU的并行计算能力，加速图像处理任务，从而实现更高的性能。通过使用OpenCV中的CUDA模块，我们可以利用GPU的并行计算能力加速图像处理和计算机视觉任务。
opencv和cuda的安装配置
serieslover4的博客
05-27 2679
opencv和cuda的安装配置工具准备安装配置安装cuda安装opencv测试和使用 工具准备 第一组： opencv3.4.5和cuda10.1，win7环境； 第二组： opencv3.4.5和cuda9.2，win10环境； 下面主要使用第一组配置 安装配置 安装cuda 1、解压并开始安装cuda 2、安装过程 安装过程中可能会出现驱动兼容性警告。此时可以选择继续安装，只要你...
OpenCV + CUDA + cuDNN模块编译
不会代码的小林
08-08 1476
在追求高端性能与资源优化并重的应用场景中，如边缘计算设备或资源受限的开发板上运行YOLO等复杂深度学习模型，采用C++结合OpenCV与GPU加速技术相较于传统的Python环境展现出显著优势。这种策略不仅极大地提升了执行效率，还显著降低了运行时资源消耗，是实现实时物体检测与识别的理想选择。C++运行YOLO，安装环境稍微复杂些，不像Python一键安装pip包。本案例使用了GPU加速算法运行,所以编译OpenCV的时候需要开启CUDA和DNN模块。
评论 3
添加红包

请填写红包祝福语或标题

个

红包个数最小为10个

元

红包金额最低5元

当前余额3.43元 前往充值 >
需支付：10.00元
成就一亿技术人!
领取后你会自动成为博主和红包主的粉丝 规则
hope_wisdom
发出的红包
实付元
使用余额支付
点击重新获取
扫码支付
钱包余额 0

抵扣说明：

1.余额是钱包充值的虚拟货币，按照1:1的比例进行支付金额的抵扣。
2.余额无法直接购买下载，可以购买VIP、付费专栏及课程。

余额充值