opencv3.3 CUDA 初学实例

//swap.cu  
   
   
#include "cuda_runtime.h"  
#include "device_launch_parameters.h"  
   
#include <opencv2/core/cuda_devptrs.hpp>  
using namespace cv;  
using namespace cv::gpu;  
   
//自定义内核函数  
__global__ void swap_rb_kernel(const PtrStepSz<uchar3> src,PtrStep<uchar3> 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);  
    }  
}  
   
extern "C" void swap_rb_caller(const PtrStepSz<uchar3>& src,PtrStep<uchar3> 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<<<grid,block,0,stream>>>(src,dst);  
    if(stream == 0)  
        cudaDeviceSynchronize();  
}  

//swap.cpp  
   
   
   
#include <opencv2/gpu/gpu.hpp>  
#include <opencv2/gpu/stream_accessor.hpp>  
   
   
using namespace cv;  
using namespace cv::gpu;  
   
extern "C" void swap_rb_caller(const PtrStepSz<uchar3>& src,PtrStep<uchar3> dst,cudaStream_t stream);  
   
extern "C" 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 <iostream>  
#include <opencv2/opencv.hpp>  
#include <opencv2/gpu/gpu.hpp>  
   
#pragma comment(lib,"opencv_gpu2410d.lib")  
#pragma comment(lib,"opencv_core2410d.lib")  
#pragma comment(lib,"opencv_highgui2410d.lib")  
   
using namespace cv;  
using namespace cv::gpu;  
   
extern "C" 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);  
    getchar();  
    waitKey(0);  
    return 0;  
}  