CUDA Samples 图像去噪实战KNN 与非局部均值 NLM 的自适应滤波实现解析【免费下载链接】cuda-samplesSamples for CUDA Developers which demonstrates features in CUDA Toolkit项目地址: https://gitcode.com/GitHub_Trending/cu/cuda-samples本文基于 NVIDIA CUDA Samples 仓库中的imageDenoising示例位于 cpp/2_Concepts_and_Techniques/imageDenoising系统讲解两种基于几何距离与颜色距离的自适应图像去噪技术——KNNK 近邻与非局部均值NLMNon-Local Means并深入剖析该示例在 CUDA 中对 NLM 所做的共享内存加速变体NLM2。读完本文你将掌握两类自适应去噪算法的权重公式与阈值调参逻辑、CUDA 纹理对象Texture Object与 OpenGL PBO 互操作的完整调用链以及该示例的构建、交互运行与自动化验证方式。示例概述与算法背景imageDenoising的官方描述见 README.md指出该示例演示了两种自适应图像去噪技术——KNN 与 NLM二者均基于纹素texel之间的几何距离与颜色距离来计算滤波权重。这两种技术在 DirectX SDK 中已有基于 shader 的实现而 CUDA 版本额外提供了一个利用共享内存shared memory大幅加速的 NLM 变体即 NLM2与 DirectX 的对应实现并列。也就是说这个示例的价值体现在三个层面算法层面展示了自适应滤波的核心思想——权重不再像高斯模糊那样只由像素间的空间距离决定而是同时结合像素颜色的相似度从而在平滑噪声的同时更好地保留边缘与纹理工程层面演示了 CUDA 纹理对象cudaTextureObject_t、CUDA Array、以及 CUDA 与 OpenGL PBOPixel Buffer Object互操作的标准用法优化层面通过 NLM2 展示了如何利用线程块内共享内存消除大量重复计算把一个朴素实现加速到可以与 DirectX shader 方案并列的程度。示例的输入图像为portrait_noise.bmp含噪人像位于 data/portrait_noise.bmpdata 目录下还提供了四种滤波模式的参考输出ref_passthru.ppm、ref_knn.ppm、ref_nlm.ppm、ref_nlm2.ppm用于自动化测试的基准比对。关键概念几何距离与颜色距离无论是 KNN 还是 NLM权重计算都建立在一个统一的基础上——距离度量。在 imageDenoising.cu 中定义了颜色距离函数__device__ float vecLen(float4 a, float4 b) { return ((b.x - a.x) * (b.x - a.x) (b.y - a.y) * (b.y - a.y) (b.z - a.z) * (b.z - a.z)); }它计算两个纹素在 RGB 颜色空间中的欧氏距离平方省略了通道的 alpha 部分。配合纹理坐标中的几何偏移(i, j)最终的权重统一表达为weight exp( -(colorDistance * Noise geometricDistance * INV_WINDOW_AREA) )其中colorDistance颜色距离体现颜色越相似权重越大geometricDistance纹素与窗口中心的几何距离平方如i*i j*jNoise噪声水平参数控制颜色距离对权重的贡献强度值越大颜色差异的惩罚越强INV_WINDOW_AREA窗口面积倒数用于将几何项归一化。该公式在两个 kernel 中都有直接对应实现例如 imageDenoising_knn_kernel.cuh 中的float weightIJ __expf(-(distanceIJ * Noise (i * i j * j) * INV_KNN_WINDOW_AREA));以及 imageDenoising_nlm_kernel.cuh 中的对应表达式。__expf()是 CUDA 提供的快速指数函数说明这套滤波在实时/交互场景下的性能取向。滤波配置参数从 imageDenoising.h 说起所有算法的核心可调参数集中定义在 imageDenoising.h是理解整个示例的关键入口宏默认值含义KNN_WINDOW_RADIUS3KNN 滤波搜索窗口半径7×7 窗口NLM_WINDOW_RADIUS3NLM 搜索窗口半径7×7 窗口NLM_BLOCK_RADIUS3NLM 相似度块patch半径7×7 块KNN_WINDOW_AREA49KNN 窗口面积(2r1)²NLM_WINDOW_AREA49NLM 窗口面积(2r1)²INV_KNN_WINDOW_AREA1/49KNN 几何项归一化系数INV_NLM_WINDOW_AREA1/49NLM 几何项归一化系数KNN_WEIGHT_THRESHOLD0.02KNN 权重阈值判定有效邻居KNN_LERP_THRESHOLD0.79KNN LERP 商阈值NLM_WEIGHT_THRESHOLD0.10NLM 权重阈值NLM_LERP_THRESHOLD0.10NLM LERP 商阈值BLOCKDIM_X / BLOCKDIM_Y8 / 8线程块维度64 线程/块权重阈值与 LERP 自适应的作用代码中每个 kernel 都维护一个fCount计数器遍历窗口时若某邻居的权重超过*_WEIGHT_THRESHOLD则累加INV_WINDOW_AREA归一化计数范围约 [0,1]。之后依据计数与*_LERP_THRESHOLD的关系决定滤波结果的融合比例float lerpQ (fCount KNN_LERP_THRESHOLD) ? lerpC : 1.0f - lerpC; clr.x lerpf(clr.x, clr00.x, lerpQ); // 在滤波结果与原像素颜色之间线性插值这段逻辑在 imageDenoising_knn_kernel.cuh 中体现得最清楚当窗口内有效相似邻居足够多时说明该区域大概率是平滑区域可以更大程度相信滤波结果反之例如位于边缘或纹理区则更多保留原始像素值从而自适应地避免过度模糊边缘。lerpC是运行时交互参数默认 0.2见下文键盘操作lerpf为设备端线性插值函数定义在 imageDenoising.cu。三大滤波 Kernel 的实现解析四个 kernel 源文件以#include方式被 imageDenoising.cu 统一引入每个文件都提供正式滤波与诊断模式diag两个版本并导出对应的cuda_*启动封装函数。PassthroughCopykernelimageDenoising_copy_kernel.cuh 定义Copykernel仅做纹理到全局内存的直通拷贝作为对比基准对应交互界面中的模式 1 Passthrough。它演示了最基本的纹理采样写法线程坐标加 0.5 偏移以对齐纹素中心通过tex2Dfloat4(texImage, x, y)采样。KNN kernel颜色近邻加权平均imageDenoising_knn_kernel.cuh 定义KNNkernel核心流程以当前纹素为中心遍历 7×7半径 3窗口内的邻居计算每个邻居与中心的颜色距离vecLen按统一权重公式得到weightIJ以权重累加颜色与权重和最后归一化依据fCount与KNN_LERP_THRESHOLD的关系将滤波结果与原始像素做 LERP 融合。其启动封装cuda_KNN采用 8×8 线程块、iDivUp向上取整划分网格该辅助函数定义在 imageDenoising.cu边界由 kernel 内if (ix imageW iy imageH)处理。NLM kernel非局部均值imageDenoising_nlm_kernel.cuh 定义NLMkernel与 KNN 的关键区别在于颜色距离不是点对点比较而是块对块比较for (float n -NLM_BLOCK_RADIUS; n NLM_BLOCK_RADIUS; n) for (float m -NLM_BLOCK_RADIUS; m NLM_BLOCK_RADIUS; m) weightIJ vecLen(tex2Dfloat4(texImage, x j m, y i n), tex2Dfloat4(texImage, x m, y n));也就是说判断邻居(xj, yi)与中心(x, y)是否相似不是看单个像素而是看两者周围 7×7半径 3邻域块的累积颜色距离。这是非局部均值算法的精髓利用图像中的自相似性重复纹理、平坦区域让远距离但结构相似的像素也能参与加权从而在去除噪声的同时保留细节。注意朴素 NLM 的计算代价很高对窗口内每个邻居都要再计算一个块距离且相邻线程的块距离计算大量重叠。这正是 NLM2 要解决的性能瓶颈。NLM2 kernel基于共享内存的加速变体imageDenoising_nlm2_kernel.cuh 定义NLM2kernel其文件头部的 8×8 网格注释清晰地说明了设计思路8×8 线程块内的每个线程只计算一个块距离ColorDistance结果存入跨线程共享的权重数组fWeights[BLOCKDIM_X * BLOCKDIM_Y]线程同步cg::sync(cta)使用 cooperative_groups 提供的线程块同步随后块内所有像素复用这批权重将其视为常数直接累加加权颜色。cg::thread_block cta cg::this_thread_block(); __shared__ float fWeights[BLOCKDIM_X * BLOCKDIM_Y]; // ... 每个线程只算一个 weight写入 fWeights ... cg::sync(cta); // ... 复用共享权重做加权平均 ...这样原本每个像素都需要计算的(2r1)² × (2r1)²次纹理采样被压缩为每个像素一次距离计算 共享内存读取纹理采样总量大幅下降。这正是 README 中所说massively speeded up variation, taking advantage of shared memory利用共享内存大幅加速的变体的工程落地也是与 DirectX shader 实现并列的性能关键。诊断模式diag平滑/边缘区域可视化每个 kernel 文件都附带*diag版本例如KNNdiag、NLMdiag、NLM2diag。它们只计算fCount并据此输出纯色标记当fCount LERP_THRESHOLD时输出红色make_color(lerpQ, 0, (1-lerpQ), 0)否则输出蓝色。运行时按下*键即可在正常模式与诊断模式间切换直观地高亮显示哪些区域被判为平滑偏向滤波、哪些区域被判为边缘/纹理偏向保留原图。数据通路纹理对象与 CUDA-OpenGL 互操作纹理对象创建imageDenoising.cu 中的CUDA_MallocArray完成了从主机图像到设备纹理的完整准备cudaMallocArray分配 CUDA Arrayuchar4通道格式cudaMemcpy2DToArray将主机图像拷入填充cudaResourceDesc资源类型为cudaResourceTypeArray与cudaTextureDesccudaCreateTextureObject创建纹理对象关键配置包括filterMode cudaFilterModeLinear线性滤波软性平滑采样addressMode[0/1] cudaAddressModeWrap坐标环绕寻址readMode cudaReadModeNormalizedFloat8 位整数被归一化为 [0,1] 浮点读取——这解释了为什么权重公式与颜色插值都以浮点进行。CUDA_FreeArray则负责cudaFreeArray释放。头文件中以extern C导出这两个函数与texImage全局纹理对象见 imageDenoising.h。OpenGL PBO 互操作imageDenoisingGL.cpp 中的initOpenGLBuffers展示了 CUDA-OpenGL 互操作的标准流程创建 OpenGL 纹理gl_Tex与 PBOgl_PBOGL_PIXEL_UNPACK_BUFFER_ARB通过cudaGraphicsGLRegisterBuffer(cuda_pbo_resource, gl_PBO, cudaGraphicsMapFlagsWriteDiscard)将 PBO 注册为 CUDA 图形资源代码注释说明旧的cudaGLRegisterBufferObject已弃用应使用新的 Graphics Resource API与 README 的 CUDA API 列表一致每帧渲染时在displayFunc中依次调用cudaGraphicsMapResources、cudaGraphicsResourceGetMappedPointer取得设备指针执行滤波 kernel再cudaGraphicsUnmapResources解锁见 imageDenoisingGL.cpp之后 OpenGL 将该 PBO 内容通过glTexSubImage2D上传为纹理并绘制全屏三角形完成显示。主循环由 GLUT 驱动timerEvent以 10msREFRESH_DELAY间隔触发重绘窗口标题实时显示当前滤镜模式与 FPS。绘制要求 OpenGL 1.5 及以上并需要GL_ARB_vertex_buffer_object与GL_ARB_pixel_buffer_object扩展见initGL中的检查逻辑。交互运行键盘控制与命令行参数交互模式下程序加载portrait_noise.bmp后进入 GLUT 主循环入口在 imageDenoisingGL.cpp支持以下按键keyboard回调按键功能1Passthrough直通显示原始带噪图像2KNN 滤波3NLM 滤波4Quick NLMNLM2共享内存加速版*切换诊断模式平滑/边缘区域红蓝高亮n/N减小 / 增大噪声水平步长 0.025l/L减小 / 增大 LERP 商lerpC范围 [0,1]步长 0.025f/F打印帧率?打印当前lerpC、knnNoise、nlmNoiseq/Esc退出运行参数的默认值定义在 imageDenoisingGL.cppknnNoise 0.32、nlmNoise 1.45、lerpC 0.2、noiseStep lerpStep 0.025。注意在调用 kernel 时代码传入的是1.0f / (noise * noise)作为Noise参数见runImageFilters即噪声水平以标准差形式交互、以方差倒数形式参与权重公式。无界面自动化测试模式示例还内置了自动化验证路径。启动时若带--file参数imageDenoisingGL.cpp则跳过 OpenGL 交互直接执行runAutoTestimageDenoising --fileoutput.ppm [--kernel0|1|2|3]--file指定输出 PPM 文件的保存路径sdkSavePPM4ub写出--kernel选择滤波模式0Passthrough1KNN2NLM3NLM2默认1。该模式在无显示器环境下也可运行main中 Linux 下会setenv(DISPLAY, :0, 0)并可通过输出文件与 data 目录中的ref_*.ppm参考图像比对验证正确性。此外交互模式下还支持-devicen指定 CUDA 设备无 OpenGL 时。构建与依赖构建该示例需要 X11、OpenGL、Freeglut、GLEW 等依赖具体说明见仓库根 README.md 中对应章节#x11、#opengl、#freeglut、#glew锚点。CMakeLists.txt 展示了其构建逻辑find_package(CUDAToolkit REQUIRED)、find_package(OpenGL)、find_package(GLUT)仅当 OpenGL 与 GLUT 均找到时才构建目标imageDenoising否则打印跳过提示GLUT not found - will not build sample源文件由bmploader.cppBMP 加载见 bmploader.cpp、imageDenoising.cu、imageDenoisingGL.cpp组成编译特性为cxx_std_17 cuda_std_17启用CUDA_SEPARABLE_COMPILATION构建后自动将data目录复制到输出目录保证运行时能找到portrait_noise.bmp与参考图像Windows 下额外链接freeglut.lib/glew64.lib或glew32.lib并拷贝对应 DLL。典型构建流程Linuxcd cuda-samples mkdir build cd build cmake .. -DCMAKE_CUDA_ARCHITECTURES75;80;86;89;90 make -j$(nproc)随后进入build/cpp/2_Concepts_and_Techniques/imageDenoising运行./imageDenoising即可。支持平台与涉及的 CUDA API按官方 README 声明支持的 SM 架构SM 5.0、5.2、5.3、6.0、6.1、7.0、7.2、7.5、8.0、8.6、8.7、8.9、9.0支持的操作系统Linux、Windows支持的 CPU 架构x86_64、armv7l依赖X11、OpenGL、Freeglut、GLEW详见上文构建小节前置条件安装对应平台的 CUDA Toolkit。涉及的 CUDA Runtime API 包括cudaMallocArray、cudaFreeArray、cudaMemcpy、cudaMalloc、cudaFree、cudaCreateTextureObject、cudaDeviceSynchronize、cudaGraphicsMapResources、cudaGraphicsUnmapResources、cudaGraphicsResourceGetMappedPointer、cudaGraphicsGLRegisterBuffer、cudaGraphicsUnregisterResource以及旧接口cudaGLRegisterBufferObject——完整清单见 README.md。延伸阅读关于 KNN 与 NLM 算法的更详细数学推导与实验对比示例随附了白皮书文档 doc/imageDenoising.pdf另有 Word 版 doc/imageDenoising.docdoc 目录下的NLM_sm.png、NLM_md.png、NLM_lg.png则提供了不同尺度下 NLM 处理结果的直观对照。小结imageDenoising是一个麻雀虽小、五脏俱全的 CUDA 图像处理样例算法上覆盖了 KNN 与 NLM 两类自适应滤波并给出了从朴素 NLM 到共享内存加速 NLM2 的完整演进工程上串联了 BMP 加载、CUDA Array 纹理对象、OpenGL PBO 互操作、GLUT 交互与无界面自动化验证五条技术线。理解本示例既能为自研图像去噪 kernel 提供可直接套用的权重公式与调参思路也能作为学习 CUDA-OpenGL 互操作与共享内存优化的标准范本。【免费下载链接】cuda-samplesSamples for CUDA Developers which demonstrates features in CUDA Toolkit项目地址: https://gitcode.com/GitHub_Trending/cu/cuda-samples创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考
