参考http://www.cnblogs.com/dwdxdy/p/3528711.html博客,加之以实践环节,我们可以知道有几种使用到GPU运算的方法:
1.利用OpenCV中提供的GPU模块
2.单独使用Cuda API编程
3.利用OpenCV中提供接口,并结合Cuda API编程
如果仅仅使用OpenCV中的GPU函数,就像博客(三)中演示一下,的确非常的简单而且可以得到比较理想的效果,但是缺点也是显然的,这种直接利用别人的函数是非常不灵活的。很多情况下,并行计算都会存在一个最优的问题,如何使得实际问题计算最优化,这就需要我们自己去思考,而不是仅仅调用别人的函数;
如果光单独使用CUDA编程,那这个工作量是可想而知的了,上文中有说道,GPU是处理计算量大计算繁琐复杂的部分,而其他的部分交给CPU去完成就是,那么这一部分完全可以使用直接提供的函数,例如OpenCV,所以,光使用CUDA进行编程,我也觉得没有必要;
最后,针对我所研究的课题,OpenCV结合CUDA编程是再适合不过的了。
这一次,我主要介绍在Ubuntu下如何将OpenCV与CUDA结合起来,也就是如何将.CU文件和.CPP文件结合起来得到想要的结果。这一部分我在网上查阅过大量的资料,但是资料非常的少,有的也只是很含糊的一笔带过,虽然方法非常简单,但是这里我还是详细地说明一下。
首先,我们的目的是:OpenCV+CUDA,完成对图像的灰度化
方法非常的简单,就是利用平均值算法 gray = ( r + g + b ) / 3得到对应像素点的值,那么这个计算部分我们就交给CUDA完成,最后结果我们将返还给CPU进行进一步处理以及结果显示。
先上代码:
test.cu文件
#include <opencv2/core/cuda_devptrs.hpp> #include "device_launch_parameters.h" #include "cuda.h" #include "cuda_runtime_api.h" #include "book.h" using namespace cv; using namespace cv::gpu; //自定义内核函数 __global__ void swap_rb_kernel(const PtrStepSz<uchar3> src, PtrStep<unsigned char> dst) { int y = blockIdx.y*blockDim.y+threadIdx.y; int x = blockIdx.x*blockDim.x+threadIdx.x; if(x < src.cols && y < src.rows) { dst(y, x) = (src(y, x).x + src(y, x).y + src(y, x).z) / 3; } } extern "C" void swap_rb_caller(const PtrStepSz<uchar3>& src, PtrStep<unsigned char> dst, cudaStream_t stream) { int uint = 8; dim3 grid(src.cols + uint-1/uint,src.rows + uint-1/uint); dim3 block(uint,uint); swap_rb_kernel<<<grid,block>>>(src, dst); if(stream == 0) cudaDeviceSynchronize(); }
testcpp.cpp
#include "stdio.h"
#include "opencv2/core/core.hpp"
#include "opencv2/highgui/highgui.hpp"
#include "opencv2/imgproc/imgproc.hpp"
#include "opencv2/gpu/gpu.hpp"
#include "opencv2/ml/ml.hpp"
#include "opencv2/gpu/gpu.hpp"
#include <opencv2/gpu/stream_accessor.hpp>
#include "cuda.h"
#include "cuda_runtime_api.h"
#include <cmath>
#include <iostream>
#ifdef _DEBUG
#pragma comment(lib, "opencv_core249d.lib")
#pragma comment(lib, "opencv_imgproc249d.lib")
#pragma comment(lib, "opencv_highgui249d.lib")
#else
#pragma comment(lib, "opencv_core249.lib")
#pragma comment(lib, "opencv_imgproc249.lib")
#pragma comment(lib, "opencv_highgui249.lib")
#endif // DEBUG
using namespace cv;
using namespace cv::gpu;
extern "C" 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(), CV_8UC1);
cudaStream_t s = StreamAccessor::getStream(stream);
swap_rb_caller(src, dst, s);
//cuResize(src, src.getWidth(), src.getHeight(), dst, dst.getWidth(), dst.getHeight())
}
int main()
{
cudaSetDevice(0);
cudaFree(0);
printf("hello\n");
Mat image = imread("test.png");
imshow("src",image);
GpuMat gpuMat,output;
output.create(image.size(), CV_8UC1);
Mat result;
result.create(image.size(), CV_8UC1);
double start = (double)getTickCount();
gpuMat.upload(image);
swap_rb(gpuMat,output);
output.download(result);
double t = ((double)getTickCount() - start)/getTickFrequency();
printf("gpu time is : %d ms\n", (int)(1000 * t));
Mat result2;
start = (double)getTickCount();
cvtColor(image, result2, COLOR_BGR2GRAY);
t = ((double)getTickCount() - start)/getTickFrequency();
printf("cpu time is : %d ms\n", (int)(1000 * t));
imshow("gpu", result);
imshow("cpu", result2);
waitKey(0);
return 0;
}
这么一看,这个问题实际上就变成了CPP中调用CU中定义的函数
这里我相信大家已经看到了, 从CPU传送数据到GPU实际上就是GpuMat->PtrStepSz的过程,注意,如果Mat是3通道的,就要PtrStepSz<uchar3>,一通道的则为PtrStepSz<unsigned char>。
关于PtrStepSz的具体信息,可以参考http://docs.opencv.org/2.4/modules/gpu/doc/data_structures.html
有了程序,接下来我们尝试运行它。
首先,编译.CU文件,在终端下输入命令:
nvcc test.cu -c -o test
编译成功后生成了一个test的文件,类型暂不明
终端输入命令:
ar cqs libtest5.a test
则生成了一个libtest5.a的静态库,这个库名可以自己另取
接下来我们要将库考入ubuntu默认搜素路径:
sudo cp libtest5.a /usr/lib/
于是,接下来可以编译cpp程序:
g++ testcpp.cpp -lopencv_core -lopencv_imgproc -lopencv_highgui -lopencv_calib3d -lopencv_contrib -lopencv_features2d -lopencv_flann -lopencv_gpu -lopencv_legacy -lopencv_ml -lopencv_objdetect -lopencv_photo -lopencv_stitching -lopencv_superres -lopencv_video -lopencv_videostab -L./ -ltest5 -o testcpp -lcudart -lpthread
这里一定要注意了,如果报出DSO missing from command line这个错误,第一看看相应的包是否导入/usr/iib/中,第二,如果已经导入成功了,就将-l语句写到命令的最后,编译才不会出错!!!!一定注意!
于是,就可以生成我们熟悉的可执行文件
运行得到结果:
运行成功,但是时间表明,对于灰度化处理,gpu运行时间并没有cpu快,这说明,传输消耗的时间远比运行时间长,这样使用gpu计算是划不来的。
但是本人的目的是成功运行cuda+opencv程序,至此,已经成功!