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