cuda做lut 去畸变示例
LUT 去畸变的原理可以简单理解为 “提前做好地图,用时直接查”,不需要复杂公式也能明白:
核心思路
镜头会让图像产生畸变(比如直线变弯),去畸变就是要把这些变形的图像 “修正” 回正常样子。
LUT 方法的关键是:提前算好一张 “映射表”(就像地图),记录下 “正常图像的每个点” 对应 “畸变图像里的哪个位置”。这样处理时不用每次重新计算,直接查这张表就能快速修正。
具体步骤
- 提前制作 “映射表”(LUT 表)
假设我们要生成一张 “无畸变的理想图像”,先想象出这张图上的每个像素点(比如左上角第一个点、右边第二个点等)。
对每个理想像素点,根据相机的特性(内参、畸变系数),计算出:这个点在原始的畸变图像中,对应的位置是哪里(可能不是整数,比如在第 100.3 行、200.5 列附近)。
把这些对应关系存成两张表:lutX 记录 X 方向坐标,lutY 记录 Y 方向坐标。 - 用 “映射表” 修正图像
拿到一张畸变的原始图像后,要生成无畸变图像时,就逐个处理每个像素:
对于理想图像中的每个点(u, v),去查 LUT 表,找到它在原始畸变图像中对应的位置(x, y)。
因为(x, y)可能不是整数点,所以取它周围几个真实像素的颜色,通过平均计算出这个位置的颜色(这一步叫 “插值”)。
把计算出的颜色填到理想图像的(u, v)位置,全部算完就得到了去畸变后的图像。
为什么好用?
快:制作 LUT 表只需要一次,之后处理每张图像时,不用重复计算复杂的畸变关系,直接查表就行,特别适合视频实时处理。
灵活:同一台相机的 LUT 表可以反复用,换相机了重新做一张表就行。
简单说,就是 “提前算好所有对应关系,用的时候直接按图索骥”,用空间换时间,实现高效的畸变校正。
#include <opencv2/opencv.hpp>
#include <iostream>
#include <vector>using namespace cv;
using namespace std;// 生成去畸变LUT表(CPU版本)
void generateUndistortLUT(Size correctedSize,const Mat& cameraMatrix,const Mat& distCoeffs,Mat& lutX,Mat& lutY)
{// 初始化LUT表,存储每个校正后像素对应的原始图像坐标lutX.create(correctedSize, CV_32F);lutY.create(correctedSize, CV_32F);// 提取相机内参和畸变系数double fx = cameraMatrix.at<double>(0, 0);double fy = cameraMatrix.at<double>(1, 1);double cx = cameraMatrix.at<double>(0, 2);double cy = cameraMatrix.at<double>(1, 2);double k1 = distCoeffs.at<double>(0);double k2 = distCoeffs.at<double>(1);double p1 = distCoeffs.at<double>(2);double p2 = distCoeffs.at<double>(3);double k3 = distCoeffs.at<double>(4);// 遍历校正后图像的每个像素,计算对应的原始畸变图像坐标for (int v = 0; v < correctedSize.height; ++v){for (int u = 0; u < correctedSize.width; ++u){// 1. 将校正后像素坐标(u, v)转换为归一化坐标double X = (u - cx) / fx;double Y = (v - cy) / fy;// 2. 计算径向畸变因子double r2 = X*X + Y*Y;double r4 = r2 * r2;double r6 = r4 * r2;double radialFactor = 1 + k1*r2 + k2*r4 + k3*r6;// 3. 计算畸变后的归一化坐标double X_distort = X * radialFactor + 2*p1*X*Y + p2*(r2 + 2*X*X);double Y_distort = Y * radialFactor + p1*(r2 + 2*Y*Y) + 2*p2*X*Y;// 4. 转换回原始图像像素坐标并存储到LUT表lutX.at<float>(v, u) = static_cast<float>(X_distort * fx + cx);lutY.at<float>(v, u) = static_cast<float>(Y_distort * fy + cy);}}
}// 使用LUT表进行去畸变(CPU版本,双线性插值)
Mat undistortWithLUT(const Mat& distortedImg,const Mat& lutX,const Mat& lutY)
{// 检查输入有效性CV_Assert(distortedImg.type() == CV_8UC3);CV_Assert(lutX.type() == CV_32F && lutY.type() == CV_32F);CV_Assert(lutX.size() == lutY.size());// 初始化输出图像Size correctedSize = lutX.size();Mat correctedImg(correctedSize, CV_8UC3);// 遍历校正后图像的每个像素for (int v = 0; v < correctedSize.height; ++v){for (int u = 0; u < correctedSize.width; ++u){// 1. 从LUT表获取原始图像坐标float x = lutX.at<float>(v, u);float y = lutY.at<float>(v, u);// 2. 边界检查(超出范围的像素设为黑色)if (x < 0 || x >= distortedImg.cols - 1 || y < 0 || y >= distortedImg.rows - 1){correctedImg.at<Vec3b>(v, u) = Vec3b(0, 0, 0);continue;}// 3. 双线性插值计算像素值int x0 = static_cast<int>(x); // 整数部分(左上角x)int y0 = static_cast<int>(y); // 整数部分(左上角y)float dx = x - x0; // 小数部分(x方向)float dy = y - y0; // 小数部分(y方向)// 4. 获取周围四个像素的值Vec3b p00 = distortedImg.at<Vec3b>(y0, x0); // 左上角Vec3b p01 = distortedImg.at<Vec3b>(y0, x0 + 1); // 右上角Vec3b p10 = distortedImg.at<Vec3b>(y0 + 1, x0); // 左下角Vec3b p11 = distortedImg.at<Vec3b>(y0 + 1, x0 + 1); // 右下角// 5. 计算插值结果(四个像素的加权平均)for (int c = 0; c < 3; ++c) // 分别处理B、G、R三个通道{correctedImg.at<Vec3b>(v, u)[c] = static_cast<uchar>(p00[c] * (1 - dx) * (1 - dy) +p01[c] * dx * (1 - dy) +p10[c] * (1 - dx) * dy +p11[c] * dx * dy);}}}return correctedImg;
}int main()
{// 1. 相机内参和畸变系数(可替换为实际标定结果)Mat cameraMatrix = (Mat_<double>(3, 3) << 500, 0, 320, // fx, 0, cx0, 500, 240, // 0, fy, cy0, 0, 1); // 0, 0, 1Mat distCoeffs = (Mat_<double>(5, 1) << -0.2, 0.1, 0.001, -0.002, 0.05); // k1, k2, p1, p2, k3// 2. 读取畸变图像Mat distortedImg = imread("distorted_image.jpg");if (distortedImg.empty()){cerr << "无法读取畸变图像!" << endl;return -1;}// 3. 生成LUT表(校正后图像尺寸与原始图像相同)Size correctedSize = distortedImg.size();Mat lutX, lutY;generateUndistortLUT(correctedSize, cameraMatrix, distCoeffs, lutX, lutY);cout << "LUT表生成完成!" << endl;// 4. 使用LUT表进行去畸变Mat correctedImg = undistortWithLUT(distortedImg, lutX, lutY);// 5. 保存并显示结果imwrite("corrected_image.jpg", correctedImg);imshow("原始畸变图像", distortedImg);imshow("LUT去畸变结果", correctedImg);waitKey(0);return 0;
}
cuda 加速
使用 CUDA 加速 CPU 版本的 LUT 去畸变过程,核心思路是将计算密集型的循环操作并行化,充分利用 GPU 的多线程架构。主要加速点集中在两个关键步骤:LUT 表生成和基于 LUT 的图像插值,具体思路如下:
一、整体加速思路
GPU 的优势在于并行处理大量重复计算。对于图像像素级别的操作(每个像素的计算逻辑独立),可以为每个像素分配一个线程,实现并行计算,替代 CPU 的串行循环。
二、分步加速方案
- LUT 表生成的 CUDA 加速
CPU 版本中,generateUndistortLUT函数通过双层 for 循环遍历每个像素计算映射关系,这是典型的可并行场景:
并行化策略:
为校正后图像的每个像素(u, v)分配一个 GPU 线程,同时计算其在原始畸变图像中的对应坐标(x, y)。
实现步骤:
将相机内参(fx, fy, cx, cy)和畸变系数(k1~k3, p1~p2)从 CPU 复制到 GPU 全局内存。
配置线程网格(grid)和线程块(block):通常用 2D 结构(如block(32, 32),grid尺寸根据图像分辨率计算)。
在 CUDA 核函数中,每个线程负责计算一个像素的映射关系,直接写入全局内存中的d_lutX和d_lutY(GPU 上的 LUT 表)。
计算完成后,将 GPU 上的 LUT 表复制回 CPU 内存(供后续步骤使用,或重复用于其他图像)。
2. 图像去畸变(插值)的 CUDA 加速
CPU 版本中,undistortWithLUT函数同样通过双层 for 循环遍历像素并执行双线性插值,这也是并行化的重点:
并行化策略:
为校正后图像的每个像素(u, v)分配一个 GPU 线程,根据 LUT 表查询原始图像坐标并完成插值计算。
实现步骤:
将原始畸变图像数据、LUT 表(lutX, lutY)从 CPU 复制到 GPU 全局内存。
配置与 LUT 生成相同的线程网格和线程块结构(与图像分辨率匹配)。
在 CUDA 核函数中,每个线程执行:
从 LUT 表读取当前像素(u, v)对应的原始图像坐标(x, y)。
执行边界检查,超出范围则填充黑色。
计算双线性插值所需的 4 个邻域像素坐标及权重。
从原始图像中读取邻域像素值,计算插值结果并写入输出图像。
计算完成后,将 GPU 上的校正后图像复制回 CPU 内存,用于保存或显示。
三、关键技术细节
内存优化:
原始图像和 LUT 表数据量大,需使用cudaMalloc分配全局内存,并通过cudaMemcpy在 CPU 和 GPU 间传输数据。
对于频繁访问的 LUT 表,可考虑使用纹理内存(Texture Memory)优化访问速度(尤其对 2D 数据的随机访问)。
线程配置:
采用 2D 线程块(如32×32,共 1024 个线程,符合 GPU 硬件限制)。
网格尺寸根据图像分辨率动态计算(如gridWidth = (imageWidth + 31) / 32,确保覆盖所有像素)。
同步与错误处理:
使用cudaDeviceSynchronize()确保核函数执行完成后再进行后续操作。
加入cudaGetLastError()检查核函数启动错误,便于调试。
四、加速效果
LUT 表生成:对于 1920×1080 分辨率,CPU 需执行约 200 万次串行计算,GPU 可并行处理所有像素,耗时通常降低 10~50 倍。
图像插值:同样分辨率下,双线性插值的并行计算可实现 50~100 倍加速(取决于 GPU 性能)。
整体流程:端到端处理速度提升显著,尤其适合视频流(30fps 以上)实时去畸变场景。
通过上述思路,可将原本依赖 CPU 串行循环的 LUT 去畸变过程,改造为 GPU 并行实现,充分发挥 CUDA 的计算优势,在保证校正精度的同时大幅提升处理效率。
#include <iostream>
#include <vector>// CUDA核函数:生成去畸变LUT表
__global__ void generateLUTKernel(float* __restrict__ lutX,float* __restrict__ lutY,int width,int height,float fx,float fy,float cx,float cy,float k1,float k2,float p1,float p2,float k3)
{// 计算线程对应的像素坐标int u = blockIdx.x * blockDim.x + threadIdx.x;int v = blockIdx.y * blockDim.y + threadIdx.y;// 边界检查if (u >= width || v >= height) return;// 计算LUT表索引int idx = v * width + u;// 1. 将校正后像素坐标(u, v)转换为归一化图像坐标(X, Y)float X = (u - cx) / fx;float Y = (v - cy) / fy;// 2. 计算径向畸变因子float r2 = X * X + Y * Y;float r4 = r2 * r2;float r6 = r4 * r2;float radialFactor = 1.0f + k1 * r2 + k2 * r4 + k3 * r6;// 3. 计算畸变后的归一化坐标(X_distort, Y_distort)float X_distort = X * radialFactor + 2.0f * p1 * X * Y + p2 * (r2 + 2.0f * X * X);float Y_distort = Y * radialFactor + p1 * (r2 + 2.0f * Y * Y) + 2.0f * p2 * X * Y;// 4. 将畸变后的归一化坐标转换为原始图像像素坐标lutX[idx] = X_distort * fx + cx;lutY[idx] = Y_distort * fy + cy;
}// CUDA核函数:使用LUT表进行去畸变(双线性插值)
__global__ void undistortKernel(const uchar3* __restrict__ distortedImg,const float* __restrict__ lutX,const float* __restrict__ lutY,uchar3* __restrict__ correctedImg,int width,int height,int distortedWidth,int distortedHeight)
{// 计算线程对应的像素坐标int u = blockIdx.x * blockDim.x + threadIdx.x;int v = blockIdx.y * blockDim.y + threadIdx.y;// 边界检查if (u >= width || v >= height) return;// 计算LUT表索引int idx = v * width + u;// 从LUT表获取原始图像坐标float x = lutX[idx];float y = lutY[idx];// 边界检查(超出原始图像范围的像素设为黑色)if (x < 0.0f || x >= distortedWidth - 1.0f || y < 0.0f || y >= distortedHeight - 1.0f) {correctedImg[idx] = make_uchar3(0, 0, 0);return;}// 双线性插值计算int x0 = static_cast<int>(x);int y0 = static_cast<int>(y);float dx = x - x0;float dy = y - y0;// 计算邻域像素索引int idx00 = y0 * distortedWidth + x0;int idx01 = y0 * distortedWidth + (x0 + 1);int idx10 = (y0 + 1) * distortedWidth + x0;int idx11 = (y0 + 1) * distortedWidth + (x0 + 1);// 获取邻域像素值uchar3 p00 = distortedImg[idx00];uchar3 p01 = distortedImg[idx01];uchar3 p10 = distortedImg[idx10];uchar3 p11 = distortedImg[idx11];// 双线性插值计算每个通道float w1 = (1.0f - dx) * (1.0f - dy);float w2 = dx * (1.0f - dy);float w3 = (1.0f - dx) * dy;float w4 = dx * dy;uchar3 result;result.x = static_cast<uchar>(p00.x * w1 + p01.x * w2 + p10.x * w3 + p11.x * w4);result.y = static_cast<uchar>(p00.y * w1 + p01.y * w2 + p10.y * w3 + p11.y * w4);result.z = static_cast<uchar>(p00.z * w1 + p01.z * w2 + p10.z * w3 + p11.z * w4);correctedImg[idx] = result;
}// 使用CUDA生成LUT表
void generateUndistortLUT_CUDA(cv::Size correctedSize,const cv::Mat& cameraMatrix,const cv::Mat& distCoeffs,cv::Mat& lutX,cv::Mat& lutY)
{// 初始化输出LUT表lutX.create(correctedSize, CV_32F);lutY.create(correctedSize, CV_32F);// 提取相机参数float fx = cameraMatrix.at<double>(0, 0);float fy = cameraMatrix.at<double>(1, 1);float cx = cameraMatrix.at<double>(0, 2);float cy = cameraMatrix.at<double>(1, 2);float k1 = distCoeffs.at<double>(0);float k2 = distCoeffs.at<double>(1);float p1 = distCoeffs.at<double>(2);float p2 = distCoeffs.at<double>(3);float k3 = distCoeffs.at<double>(4);// 设备内存指针float *d_lutX, *d_lutY;size_t lutSize = correctedSize.width * correctedSize.height * sizeof(float);// 分配设备内存cudaMalloc(&d_lutX, lutSize);cudaMalloc(&d_lutY, lutSize);// 配置线程块和网格大小dim3 block(32, 32);dim3 grid((correctedSize.width + block.x - 1) / block.x,(correctedSize.height + block.y - 1) / block.y);// 启动LUT生成核函数generateLUTKernel<<<grid, block>>>(d_lutX, d_lutY,correctedSize.width, correctedSize.height,fx, fy, cx, cy,k1, k2, p1, p2, k3);cudaDeviceSynchronize();// 将结果复制回主机cudaMemcpy(lutX.data, d_lutX, lutSize, cudaMemcpyDeviceToHost);cudaMemcpy(lutY.data, d_lutY, lutSize, cudaMemcpyDeviceToHost);// 释放设备内存cudaFree(d_lutX);cudaFree(d_lutY);
}// CUDA去畸变函数
cv::Mat undistortWithCuda(const cv::Mat& distortedImg,const cv::Mat& lutX,const cv::Mat& lutY)
{CV_Assert(distortedImg.type() == CV_8UC3);CV_Assert(lutX.type() == CV_32F && lutY.type() == CV_32F);CV_Assert(lutX.size() == lutY.size());cv::Size correctedSize = lutX.size();cv::Mat correctedImg(correctedSize, CV_8UC3);// 设备内存指针uchar3* d_distorted;float* d_lutX;float* d_lutY;uchar3* d_corrected;// 计算内存大小size_t imgSize = distortedImg.rows * distortedImg.cols * sizeof(uchar3);size_t lutSize = correctedSize.width * correctedSize.height * sizeof(float);// 分配设备内存cudaMalloc(&d_distorted, imgSize);cudaMalloc(&d_lutX, lutSize);cudaMalloc(&d_lutY, lutSize);cudaMalloc(&d_corrected, imgSize);// 转换OpenCV图像格式并复制到设备cudaMemcpy(d_distorted, distortedImg.data, imgSize, cudaMemcpyHostToDevice);cudaMemcpy(d_lutX, lutX.data, lutSize, cudaMemcpyHostToDevice);cudaMemcpy(d_lutY, lutY.data, lutSize, cudaMemcpyHostToDevice);// 配置线程块和网格大小dim3 block(32, 32);dim3 grid((correctedSize.width + block.x - 1) / block.x,(correctedSize.height + block.y - 1) / block.y);// 启动去畸变核函数undistortKernel<<<grid, block>>>(d_distorted,d_lutX,d_lutY,d_corrected,correctedSize.width,correctedSize.height,distortedImg.cols,distortedImg.rows);cudaDeviceSynchronize();// 将结果复制回主机cv::Mat resultBGR(correctedSize, CV_8UC3);cudaMemcpy(resultBGR.data, d_corrected, imgSize, cudaMemcpyDeviceToHost);correctedImg = resultBGR.clone();// 释放设备内存cudaFree(d_distorted);cudaFree(d_lutX);cudaFree(d_lutY);cudaFree(d_corrected);return correctedImg;
}int main() {cv::Mat cameraMatrix = (cv::Mat_<double>(3, 3) << 799.7870974695102, 0, 306.1514141859209,0, 761.1184471649692, 310.0513553191973,0, 0, 1);// 畸变系数 (k1, k2, p1, p2, k3)cv::Mat distCoeffs = (cv::Mat_<double>(5, 1) << -0.6358100452675683,0.2841255094561374,-0.05567168566375429,0.02016180203610067,-0.3773069761975876);// 2. 读取畸变图像cv::Mat distortedImg = cv::imread("distorted_image.jpg");if (distortedImg.empty()) {std::cerr << "无法读取图像!" << std::endl;return -1;}// 3. 使用CUDA生成LUT表cv::Size correctedSize = distortedImg.size();cv::Mat lutX, lutY;generateUndistortLUT_CUDA(correctedSize, cameraMatrix, distCoeffs, lutX, lutY);// 4. 使用CUDA进行去畸变cv::Mat correctedImg = undistortWithCuda(distortedImg, lutX, lutY);bool saveSuccess = cv::imwrite("undistorted_result2.jpg", correctedImg);if (saveSuccess) {std::cout << "去畸变后的图片已保存为:undistorted_result.jpg" << std::endl;} else {std::cerr << "保存图片失败!" << std::endl;return -1;}// 5. 显示结果cv::imwrite("undistorted_result.jpg", correctedImg); // 保存去畸变后的图片cv::waitKey(0);return 0;
}
实例图片:opencv
编译运行:
nvcc lut.cu -o lut pkg-config --cflags --libs opencv4