首页 > 编程语言 >利用CUDA编程实现在GPU中对图像的极坐标变换加速

利用CUDA编程实现在GPU中对图像的极坐标变换加速

时间:2024-12-31 14:54:51浏览次数:3  
标签:polar min int float 极坐标 CUDA theta GPU cv

问题来源:

1. 需要对输入图像中的一个环形区域,进行极坐标逆变换,将该环形区域转换为一张新的矩形图像

2. opencv没有直接对环形区域图像进行变换的函数,需要通过循环遍历的方式,利用polarToCart进行转换

3. 循环遍历不可避免的带来速度上的问题,尤其是图片较大时

解决思路

1:使用opencv自带的多线程parallel_for_进行加速

2:利用cuda编程在gpu中加速

3:为了便于在cuda中实现相应功能,坐标转换的相关代码根据原理手动实现

一、利用opencv的parallel_for_多线程运算

    void polar2cart_cv(cv::Mat& mat1, cv::Mat& mat2,
        const cv::Point2f& center, const float& min_radius,
        const float& max_radius, const float& min_theta, const float& max_theta)
    {
        //转换后图像的行即圆环的半径差,列数即最大半径对应的弧长
        int rows_c = static_cast<int>(std::ceil(max_radius - min_radius));
        int cols_c = static_cast<int>(std::ceil((max_theta - min_theta) * max_radius));

        mat2 = cv::Mat::zeros(rows_c, cols_c, CV_8UC1);
        // 极坐标转换的角度步长
        float rstep = 1.0;
        float thetaStep = (max_theta - min_theta) / cols_c;

        float center_polar_x = center.x;
        float center_polar_y = center.y;

        cv::parallel_for_(cv::Range(0, rows_c * cols_c), [&](const cv::Range& range) {
            for (int r = range.start; r < range.end; r++)
            {
                int row = r / cols_c;
                int col = r % cols_c;
                float temp_r = min_radius + row * rstep;

                float theta_p = min_theta + col * thetaStep;
                float sintheta = std::sin(theta_p);
                float costheta = std::cos(theta_p);

                int polar_x = static_cast<int>(center_polar_x + temp_r * sintheta);
                int polar_y = static_cast<int>(center_polar_y + temp_r * costheta);

                if (polar_x >= 0 && polar_x < mat1.cols && polar_y >= 0 && polar_y < mat1.rows)
                {
                    mat2.at<uchar>(rows_c - 1 - row, col) = mat1.at<uchar>(polar_y, polar_x);
                }
            }
            });
    }

二、利用CUDA编程在GPU中并行运算

    void polar2cart_cuda(cv::Mat& mat1, cv::Mat& mat2,
        const cv::Point2f& center, const float& min_radius,
        const float& max_radius, const float& min_theta, const float& max_theta)
    {
        int rows_c = static_cast<int>(std::ceil(max_radius - min_radius));
        int cols_c = static_cast<int>(std::ceil((max_theta - min_theta) * max_radius));

        mat2 = cv::Mat::zeros(rows_c, cols_c, CV_8UC1);
        float rstep = 1.0;
        float thetaStep = (max_theta - min_theta) / cols_c;

        float center_polar_x = center.x;
        float center_polar_y = center.y;

        uchar* deviceMat1Data;
        cudaMalloc((void**)&deviceMat1Data, mat1.rows * mat1.cols * sizeof(uchar));

        uchar* deviceMat2Data;
        cudaMalloc((void**)&deviceMat2Data, mat2.rows * mat2.cols * sizeof(uchar));

        //图像 cpu -> gpu
        cudaMemcpy(deviceMat1Data, mat1.data, mat1.rows * mat1.cols * sizeof(uchar), cudaMemcpyHostToDevice);
        cudaMemcpy(deviceMat2Data, mat2.data, mat2.rows * mat2.cols * sizeof(uchar), cudaMemcpyHostToDevice);

        int maxRows = std::max(mat1.rows, mat2.rows);
        int maxCols = std::max(mat1.cols, mat2.cols);

        dim3 blockSize(32, 32);
        dim3 gridSize((maxCols + blockSize.x - 1) / blockSize.x, (maxRows + blockSize.y - 1) / blockSize.y);

        process<<<gridSize, blockSize>>>(deviceMat1Data, mat1.rows, mat1.cols, deviceMat2Data, mat2.rows, mat2.cols,1,1, min_radius,min_theta,thetaStep,center_polar_x,center_polar_y);

        //结果图像 gpu -> cpu
        cudaMemcpy(mat2.data, deviceMat2Data, mat2.rows * mat2.cols * sizeof(uchar), cudaMemcpyDeviceToHost);

        cudaFree(deviceMat1Data);
        cudaFree(deviceMat2Data);

    }

gpu中进行极坐标逆变换的核心函数

    __global__ void process(uchar* mat1Data, int mat1Rows, int mat1Cols, uchar* mat2Data, int mat2Rows, int mat2Cols,int channels1,int channels2,float min_radius,float min_theta,float thetaStep,float center_polar_x,float center_polar_y)
    {
        int row = blockIdx.y * blockDim.y + threadIdx.y;
        int col = blockIdx.x * blockDim.x + threadIdx.x;

        if (row < mat2Rows && col < mat2Cols) {
            
            float theta_p = min_theta + col * thetaStep;
            float sintheta = std::sin(theta_p);
            float costheta = std::cos(theta_p);

            //每个极坐标系下的点,在原图(直角坐标系)下的坐标
            float temp_r = min_radius + row*1;
            int src_col = static_cast<int>(center_polar_x + temp_r * sintheta);
            int src_row = static_cast<int>(center_polar_y + temp_r * costheta);
            
            // 结果图排放时的坐标(row的定义是沿着半径增加的方向增加,与图像的索引dst_row方向相反。col则与dst_col一致)
            int dst_row = mat2Rows-1-row;
            int dst_col = col;
            if (src_col >= 0 && src_col < mat1Cols && src_row >= 0 && src_row < mat1Rows)
            {
                //每个点对应在原图(极坐标系)的索引
                int srcidex= (src_row*mat1Cols+ src_col)*channels1;
                //每个点对应在结果图(直角坐标系)的索引
                int dstIdex =  (dst_row*mat2Cols+ dst_col)*channels2;
                mat2Data[dstIdex] =  mat1Data[srcidex];     
            }
        }
    }

三、测试与结论

1.测试代码

#include <iostream>
#include <opencv2/opencv.hpp>
#include <cuda_runtime.h>
using namespace std::chrono;
#define PI 2 * std::acos(0.0)
#include "cuda_fun_api/polar2cart.h" 

int main()
{
	cv::Mat ori_image = cv::imread("..\\..\\1.png");
	cv::Mat ori_gray;
	cv::cvtColor(ori_image, ori_gray, cv::COLOR_BGR2GRAY);
	int i = 0;
	cv::Mat output_image;
	float sum_time = 0;
	int min_radius = 680;
	int max_radius = 1330;
	cv::Point2f circle_center = cv::Point2f(2102,1539);
	while(i < 20)
	{
		double t1 = cv::getTickCount();
		// opencv的多线程
		// cuda_fun_api::polar2cart_cv(ori_gray.clone(), output_image, circle_center, min_radius, max_radius, 0, 2 * PI);

		// cuda
		cuda_fun_api::polar2cart_cuda(ori_gray.clone(), output_image, circle_center, min_radius, max_radius, 0, 2 * PI);
		double t2 = cv::getTickCount();
		//剔除第一张的运行时间(在gpu开辟内存会比较耗时)
		std::cout <<  (t2 - t1) * 1000 / cv::getTickFrequency() << std::endl;
		if (i > 0)
			sum_time += (t2 - t1) * 1000 / cv::getTickFrequency();
		i++;
	}
	std::cout << "average:" <<sum_time/19 << std::endl;

	// cv::namedWindow("ori_image",0);
	// cv::circle(ori_image,circle_center,min_radius,cv::Scalar(0,0,255),10);
	// cv::circle(ori_image,circle_center,max_radius,cv::Scalar(0,0,255),10);
	// cv::resizeWindow("ori_image", ori_image.cols/5, ori_image.rows/5);
	// cv::imshow("ori_image", ori_image);

	// cv::namedWindow("output_image",0);
	// cv::resizeWindow("output_image", output_image.cols/5, output_image.rows/5);
	// cv::imshow("output_image", output_image);
	// cv::waitKey(0);

	return 0;
}

2.测试结果

对如下图像的环境区域进行极坐标转换,图像大小为1200万像素,循环测试20次。

原图:

极坐标逆变换后:

parallel_for_方式:      CUDA方式:

    

3.结论

CUDA的方式会在第一张耗时较久(GPU中开辟内存),但整体的运行效率远高于parallel_for_方式。

CUDA方式所能提升的效率和图像大小有关,图像越大,计算量越大,提升的效率越高。

四、工程配置

本文将cuda相关的函数封装到cuda_function_api库中,与主程序解耦

find_package(CUDA)
find_package(cuBLAS)

set(CUDA_SOURCE_FILES src/cuda_fun_api/polar2cart.cu)
cuda_add_library(cuda_function_api ${CUDA_SOURCE_FILES} STATIC)
target_link_libraries(cuda_function_api  ${OpenCV_LIBS} ${CUBLAS_LIBRARIES})
FILE(GLOB SOURCES src/main.cpp)

ADD_EXECUTABLE(main ${SOURCES})
TARGET_LINK_LIBRARIES(main ${OpenCV_LIBS} cuda_function_api)
target_compile_options(main PRIVATE "/MT")

五、 源码及示例:https://github.com/ALongSuper/polar_coordinate_trans.git

标签:polar,min,int,float,极坐标,CUDA,theta,GPU,cv
From: https://blog.csdn.net/Alongsuper/article/details/144844567

相关文章

  • GPU编程最佳语言
    GPU编程最佳语言‌GPU编程的最佳语言选择取决于具体的应用场景和开发者的需求。以下是几种常用的GPU编程语言及其优缺点‌:‌CUDA‌:‌优点‌:CUDA是NVIDIA推出的并行计算平台和编程模型,基于C++,提供了丰富的库和工具,适用于需要直接访问GPU硬件的高性能计算任务。CUDA具有较低的......
  • 必要性论证:将FPGA深入应用于基于CPU、CPU+GPU的人形机器人控制系统
    目录:0前言1需求侧的基本事实1.1实用化的人形机器人的控制系统必须实现感算控一体1.2感算控环路必须具备强实时性(低延迟量+低延迟抖动量)1.3感知环节必须以高帧率+高分辨率、在多个位置+多个方向并行采集人形机器人本体、环境、任务对象的多种信息1.4强实时性的感算......
  • 高性能计算-GPU编程模型(21)
    1.GPU的内存模型GPU编程数据需要从CPU主存拷贝到GPU全局存储器,所有线程共享全局存储。开辟的全局存储器空间指针在CPU代码中不能解引用使用,应在计算完结果后再拷贝回CPU主存空间。线程块内共享存储。(1)线程私有的存储有寄存器、本地内存(2)线程块内有块内线程共享的共享内......
  • 【CUDA】cuDNN:加速深度学习的核心库
    【CUDA】cuDNN:加速深度学习的核心库1.什么是cuDNN?cuDNN(CUDADeepNeuralNetworklibrary)是NVIDIA提供的一个高性能GPU加速库,专为深度学习任务中常见的操作进行了高度优化。它不仅提供了单操作的高效实现,还支持多操作融合(fusion),旨在最大化地利用NVIDIAGPU的计算......
  • 十亿行数据挑战:CUDA申请出战(从17分钟到17秒)
    文章结尾有最新热度的文章,感兴趣的可以去看看。本文是经过严格查阅相关权威文献和资料,形成的专业的可靠的内容。全文数据都有据可依,可回溯。特别申明:数据和资料已获得授权。本文内容,不涉及任何偏颇观点,用中立态度客观事实描述事情本身导读在我学习CUDA的过程中,我决定用......
  • 欧拉系统安装GPU驱动
    安装NVIDIADriver进入英伟达官网下载页面按照以上方式选择即可得到>535.113.01版本的驱动,可以实现多卡推理,小于这个版本会导致多卡训练以及推理报错虽然最新版本为550.54.15,但是535版本更加稳定,并且pytorch目前只支持到12.1,而在CUDAToolkit选择栏中没有这个版本,所以选择12.2......
  • GPU gdm /etc/X11/xorg.conf
    【本文适用环境:Redhat或CentOS】前提:nvidia-smi能正常读取GPU卡信息关闭gdm[root@host~]#systemctlstopgdm查询系统下是否存在/etc/X11/xorg.conf文件,如果不存在则执行下述步骤生成配置文件[root@host~]#nvidia-xconfig--query-gpu-infoNumberofGPUs:1GP......
  • 绕过CPU:英伟达与IBM致力推动GPU直连SSD以大幅提升性能
    绕过CPU:英伟达与IBM致力推动GPU直连SSD以大幅提升性能|Id|Title|DateAdded|SourceUrl|PostType|Body|BlogId|Description|DateUpdated|IsMarkdown|EntryName|CreatedTime|IsActive|AutoDesc|AccessPermission||-------------|-------------|---......
  • python网络编程之http longpull
    服务端:fromflaskimportFlask,request,jsonifyimporttimeapp=Flask(__name__)@app.route('/stream',methods=['GET'])defpoll():#假设这里有一个方法来检查是否有新数据#为了示例,我们简单地模拟等待数据time.sleep(5)#模拟处理时间或等待......
  • 在 K8S 中创建 Pod 是如何使用到 GPU 的: nvidia device plugin 源码分析
    本文主要分析了在K8s中创建一个Pod并申请GPU资源,最终该Pod时怎么能够使用GPU的,具体的实现原理,以及deviceplugin、nvidia-container-toolkit相关源码分析。1.概述在两篇文章中分别分享了在不同环境如何使用GPU,以及在k8s中使用GPUOperator来加速部署。在......