CUDA跟OpenCV的混合编程，注意OpenCV需要重新编译

//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;
}