第OpenCV圖像輪廓提取的實現(xiàn)2.f/159.f,4.f/159.f,5.f/159.f,4.f/159.f,2.f/159.f,

4.f/159.f,9.f/159.f,12.f/159.f,9.f/159.f,4.f/159.f,

5.f/159.f,12.f/159.f,15.f/159.f,12.f/159.f,5.f/159.f,

4.f/159.f,9.f/159.f,12.f/159.f,9.f/159.f,4.f/159.f,

2.f/159.f,4.f/159.f,5.f/159.f,4.f/159.f,2.f/159.f,

cudaMemcpyToSymbol(convolutionKernelStore,gaussianKernel5x5,sizeof(gaussianKernel5x5),0);

constunsignedintgaussianKernel5x5Offset=0;

//Sobelgradientkernels

constfloatsobelGradientX[9]=

-1.f,0.f,1.f,

-2.f,0.f,2.f,

-1.f,0.f,1.f,

constfloatsobelGradientY[9]=

1.f,2.f,1.f,

0.f,0.f,0.f,

-1.f,-2.f,-1.f,

cudaMemcpyToSymbol(convolutionKernelStore,sobelGradientX,sizeof(sobelGradientX),sizeof(gaussianKernel5x5));

cudaMemcpyToSymbol(convolutionKernelStore,sobelGradientY,sizeof(sobelGradientY),sizeof(gaussianKernel5x5)+sizeof(sobelGradientX));

constunsignedintsobelGradientXOffset=sizeof(gaussianKernel5x5)/sizeof(float);

constunsignedintsobelGradientYOffset=sizeof(sobelGradientX)/sizeof(float)+sobelGradientXOffset;

//CreaCPU/GPUimagenescompartidos

cameraframe;

unsignedchar*sourceDataDevice,*blurredDataDevice,*edgesDataDevice;

cv::Matsource(frame.size(),CV_8U,createImageBuffer(frame.size().width*frame.size().height,sourceDataDevice));

cv::Matblurred(frame.size(),CV_8U,createImageBuffer(frame.size().width*frame.size().height,blurredDataDevice));

cv::Matedges(frame.size(),CV_8U,createImageBuffer(frame.size().width*frame.size().height,edgesDataDevice));

//Crea2imagenestemporales(sobelgradients)

unsignedchar*deviceGradientX,*deviceGradientY;

cudaMalloc(deviceGradientX,frame.size().width*frame.size().height);

cudaMalloc(deviceGradientY,frame.size().width*frame.size().height);

//Loopwhilecapturaimagenes

while(1)

//Capturalaimagenenescaladegrises

cameraframe;

cvtColor(frame,source,COLOR_RGB2GRAY);

_sleep(1);

//Grabaeltiempoquedemoraelproceso

cudaEventRecord(start);

//convolutionkernelparametros

dim3cblocks(frame.size().width/16,frame.size().height/16);

dim3cthreads(16,16);

//pythagorankernelparametros

dim3pblocks(frame.size().width*frame.size().height/256);

dim3pthreads(256,1);

//gaussianblur(firstkernelinstore@0)

convolvecblocks,cthreads(sourceDataDevice,frame.size().width,frame.size().height,0,0,gaussianKernel5x5Offset,5,5,blurredDataDevice);

//sobelgradientconvolutions(xypaddingisnow2becausethereisaborderof2arounda5x5gaussianfilteredimage)

convolvecblocks,cthreads(blurredDataDevice,frame.size().width,frame.size().height,2,2,sobelGradientXOffset,3,3,deviceGradientX);

convolvecblocks,cthreads(blurredDataDevice,frame.size().width,frame.size().height,2,2,sobelGradientYOffset,3,3,deviceGradientY);

pythagoraspblocks,pthreads(deviceGradientX,deviceGradientY,edgesDataDevice);

cudaThreadSynchronize();

cudaEventRecord(stop);

//Muestratiempodeejecucion

floatms=0.0f;

cudaEventSynchronize(stop);

cudaEventElapsedTime(ms,start,stop);

std::cout"ElapsedGPUtime:"ms"milliseconds"std::endl;

//Muestraresultados

imshow("Source",frame);

imshow("Greyscale",source);

imshow("Blurred",blurred);

imshow("Sobel",edges);

//Spin

if(cv::waitKey(1)==27)break;

//Exit

cudaFreeHost(source.data);

cudaFreeHo