OpenCV和Cuda结合编程
创始人
2024-03-30 00:35:42
0

一、利用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<<>>(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文件主程序,完成数据的输入、业务的处理和数据的输出。

总结

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

相关内容

热门资讯

读播音主持专业好吗 读播音主持专业好吗挺好滴 我是哪赚经验滴
水浒传每一回情节概括 水浒传每一回情节概括这个……我觉得名著们的每一回的标题都是概括了
男人心理不成熟的表现有哪些,如... 男人心理不成熟的表现有哪些,如何去改善?爱计较,考虑事情不周全,没有责任心等等都是不成熟的表现。
来的快去的也快,爱情里到底应该... 来的快去的也快,爱情里到底应该如何掌控节奏?一定要循序渐进,不断发掘对方身上的闪光点,也要慢慢去展露...
胳膊痛,电吹风一吹就不痛了,是... 胳膊痛,电吹风一吹就不痛了,是怎么了啊受风了,着凉了把,拔罐子就能好点,受风寒了,建议贴点祛风寒、舒...
青鱼最爱吃螺蛳了,螺蛳壳很坚硬... 青鱼最爱吃螺蛳了,螺蛳壳很坚硬青鱼是怎么吃的?青鱼非常巧妙的用它的嘴巴把螺蛳肉用力的吸出来,吃掉。玉...
心情低落,用英文怎么说? 心情低落,用英文怎么说?feel downhehe英文应该这样说: I am in low sp...
抗日英雄谱丨14岁少年英雄王璞... 转自:河北新闻网1943年5月7日日本侵略者在河北省顺平县野场村残忍杀害了118名村民制造了“野场惨...
亮马河畔,她们为水上游船拧紧“... 转自:北京日报客户端7月5日,湖南省郴州资兴市东江湖白廊景区发生一起水上交通船因突发11级大风导致侧...
第七大道(00797.HK)拟... 格隆汇7月6日丨第七大道(00797.HK)公布,于2025年7月6日,卖方(为公司的间接全资附属公...