用CUDA 实现图像remap

图像remap在图像和视频帧上经常用于仿射投影变换，畸变矫正，图像拼接

下面是对应remap核函数的实现：remap就是把目标图像根据map里的对应坐标里的像素填补到结果图像里。

__global__ void remap(uchar* img_data,uchar *img_res,float *crood_x,float *crood_y, int height,int width)
{    
    int bi = blockIdx.x;
    int tix = threadIdx.x;
    int idx = bi*blockDim.x+tix;
    float x = crood_x[idx];
    float y = crood_y[idx];
    int idx_res = (int)y*width+(int)x; 
    if (idx<width*height)
    {
        img_res[idx*3+0] = img_data[idx_res*3+0];
        img_res[idx*3+1] = img_data[idx_res*3+1];
        img_res[idx*3+2] = img_data[idx_res*3+2];
    }
}

 函数调用过程：载入图像和坐标，准备remap核函数的数据和显存，执行函数后把结果拷贝回内存中。

int main()
{
    cv::FileStorage fs2("map_0.yml", cv::FileStorage::READ); //yml保存了x，y映射坐标
    cv::Mat img = cv::imread("before.jpg");
    
    cv::Mat matx, maty;  
    fs2["matx"] >> matx;  
    fs2["maty"] >> maty;  

    float* crood_x = matx.ptr<float>(0);
    float* crood_y = maty.ptr<float>(0);
    uchar* img_data = img.ptr<uchar>(0);

    uchar* imgGpu;
	cudaMalloc((void**)&imgGpu, 1920*1080*3 * sizeof(uchar));
    cudaMemcpy(imgGpu, img_data,  1920*1080*3 * sizeof(uchar), cudaMemcpyHostToDevice);

    float* crood_xGpu;
	cudaMalloc((void**)&crood_xGpu, 1920*1080* sizeof(float));
	cudaMemcpy(crood_xGpu, crood_x,  1920*1080* sizeof(float), cudaMemcpyHostToDevice);

    float* crood_yGpu;
	cudaMalloc((void**)&crood_yGpu, 1920*1080* sizeof(float));
	cudaMemcpy(crood_yGpu, crood_y,  1920*1080* sizeof(float), cudaMemcpyHostToDevice);

    uchar *img_resGpu;
    uchar img_res[1920*1080*3];
    cudaMalloc((void**)&img_resGpu, 1920*1080*3* sizeof(float));
    
    // 执行100次，看看速度 
    for (int i=0;i<100;++i)
    {
        double time0 = static_cast<double>(cv::getTickCount());	
        cudaMemcpy(imgGpu, img_data,  1920*1080*3 * sizeof(uchar), cudaMemcpyHostToDevice);
        // 2025*1024==1920*1080*3，显卡block线程数有限制最多1024
        remap<<<2025,1024>>>(imgGpu,img_resGpu,crood_xGpu,crood_yGpu,1080,1920);
        cudaDeviceSynchronize();
        cudaMemcpy(img_res, img_resGpu,  1920*1080* 3*sizeof(uchar), cudaMemcpyDeviceToHost);
        time0 = ((double)cv::getTickCount() - time0) / cv::getTickFrequency()*1000;
        printf("remap time:%f\n", time0);
    }
    cv::Mat src(1080, 1920, CV_8UC3, img_res);
    cudaDeviceReset();
    cv::imwrite("after_remap.jpg",src);
    return 0;
}

 remap前

 remap后