__global__ void swap_rb_kernel(const uchar3* src,uchar3* dst,int width,int height)
{int x = threadIdx.x + blockIdx.x * blockDim.x;int y = threadIdx.x + blockIdx.y * blockDim.y;if(x < width && y < height){uchar3 v = src[y * width + x];dst[y * width + x].x = v.z;dst[y * width + x].y = v.y;dst[y * width + x].z = v.x;}
}void swap_rb_caller(const uchar3* src,uchar3* dst,int width,int height)
{dim3 block(32,8);dim3 grid((width + block.x - 1)/block.x,(height + block.y - 1)/block.y);swap_rb_kernel<<<grid,block,0>>>(src,dst,width,height);cudaThreadSynchronize();
}int main()
{Mat image = imread("lena.jpg");imshow("src",image);size_t memSize = image.cols*image.rows*sizeof(uchar3);uchar3* d_src = NULL;uchar3* d_dst = NULL;CUDA_SAFE_CALL(cudaMalloc((void**)&d_src,memSize));CUDA_SAFE_CALL(cudaMalloc((void**)&d_dst,memSize));CUDA_SAFE_CALL(cudaMempcy(d_src,image.data,memSize,cudaMemcpyHostToDevice));swap_rb_caller(d_src,d_dst,image.cols,image.rows);CUDA_SAFE_CALL(cudaMempcy(image.data,d_dst,memSize,cudaMemcpyDeviceToHost));imshow("gpu",image);waitKey(0);CUDA_SAFE_CALL(cudaFree(d_src));CUDA_SAFE_CALL(cudaFree(d_dst));return 0;
}

//swap_rb.cu
#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);}
}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_rb.cpp
#include <opencv2/gpu/gpu.hpp>
#include <opencv2/gpu/stream_accessor.hpp>
using namespace cv;
using namespace cv::gpu;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(),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>
using namespace cv;
using namespace cv::gpu;
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);waitKey(0);return 0;
}