CUDA 双边滤波实战：基于 cuda-samples bilateralFilter 的边缘保持图像平滑与 OpenGL 交互实现

CUDA 双边滤波实战：基于 cuda-samples bilateralFilter 的边缘保持图像平滑与 OpenGL 交互实现

【免费下载链接】cuda-samplesSamples for CUDA Developers which demonstrates features in CUDA Toolkit项目地址: https://gitcode.com/GitHub_Trending/cu/cuda-samples

导读

本文深入剖析 NVIDIA cuda-samples 仓库中的 bilateralFilter 示例（路径：cpp/5_Domain_Specific/bilateralFilter），讲解如何在 CUDA 上实现边缘保持（edge-preserving）的非线性平滑滤波器——双边滤波，并通过 OpenGL 实时渲染结果。读完本文，你将掌握双边滤波的数学原理与三个核心参数（高斯标准差、欧氏距离阈值、迭代次数）对图像的差异化影响，理解 CUDA 纹理对象、常量内存、OpenGL PBO 互操作等底层实现细节，并能独立构建、运行该示例并完成 CPU/GPU 结果自动比对验证。

一、示例概述：双边滤波是什么

bilateralFilter 是 cuda-samples 中一个经典的图像处理示例，它使用 CUDA 实现边缘保持的非线性平滑滤波，并借助 OpenGL 进行渲染展示，可用于图像恢复（image recovery）与去噪（denoising）场景。

与普通高斯滤波只考虑像素之间的空间距离不同，双边滤波在计算每个像素的输出时，同时考虑两个权重因素：

• 空间距离（spatial distance / domain filter）：中心像素与其邻域像素的几何距离，决定"离得近的像素贡献大"；
• 颜色距离（color distance / range filter）：中心像素与其邻域像素在颜色空间中的欧氏距离，决定"颜色相近的像素贡献大"。

正是这种"双重加权"机制，使得双边滤波能够在平滑噪声、去除纹理的同时，保留物体边缘的清晰轮廓。算法出处为经典论文：C. Tomasi, R. Manduchi,Bilateral Filtering for Gray and Color Images, ICCV 1998。

从仓库根目录的 cpp/5_Domain_Specific/README.md 可以看到，该示例被归类于"领域特定（Domain Specific）"章节，与 recursiveGaussian（高斯模糊）、SobelFilter（Sobel 边缘检测）等同属图像处理类示例，但 bilateralFilter 的核心亮点在于"平滑与保边兼顾"以及 CUDA 与 OpenGL 的图形互操作。

二、核心参数与视觉效果：三个旋钮如何控制滤波结果

在 bilateralFilter.cpp 中，示例定义了三个核心参数及默认值：

2.1 欧氏距离阈值（euclidean delta）——保边与平滑的平衡杆

源代码头注释（bilateralFilter.cpp）对参数效果给出了精确描述：

• euclidean delta 增大：大部分细腻纹理被滤除，但所有轮廓仍然与原图一样锐利；
• euclidean delta 趋近于无穷大：颜色域权重退化为常数，滤波器退化为普通高斯滤波器；
• gaussian delta 增大：细腻纹理被模糊得更多（空间域影响增强）；
• 多次迭代：会显著"压平"图像中的颜色，但不模糊边缘，从而产生卡通效果（cartoon effect）。

在 bilateralFilter.cpp 中，示例还通过宏限定了参数的边界：

#define MIN_EUCLIDEAN_D 0.01f #define MAX_EUCLIDEAN_D 5.f #define MAX_FILTER_RADIUS 25

2.2 参数调优的交互方式

默认的 OpenGL 交互模式下，程序注册了键盘回调（keyboard 函数），运行时可以实时调节参数并立即看到效果：

三、算法原理与 CUDA 实现细节

3.1 权重公式

核心 kernel 位于 bilateral_kernel.cu。其权重由两部分乘积构成：

// 颜色域权重（range factor）：基于欧氏距离的高斯衰减 __device__ float euclideanLen(float4 a, float4 b, float d) { float mod = (b.x - a.x) * (b.x - a.x) + (b.y - a.y) * (b.y - a.y) + (b.z - a.z) * (b.z - a.z); return __expf(-mod / (2.f * d * d)); }

而每个输出像素的计算过程为：遍历(2r+1) × (2r+1)的邻域，对每个邻域像素计算factor = 高斯空间权重 × 颜色欧氏权重，再对邻域像素做加权平均后归一化（见 d_bilateral_filter kernel）：

float4 center = tex2D<float4>(rgbaTex, x, y); for (int i = -r; i <= r; i++) { for (int j = -r; j <= r; j++) { float4 curPix = tex2D<float4>(rgbaTex, x + j, y + i); factor = cGaussian[i + r] * cGaussian[j + r] * // domain factor euclideanLen(curPix, center, e_d); // range factor t += factor * curPix; sum += factor; } } od[y * w + x] = rgbaFloatToInt(t / sum);

其中cGaussian是声明在设备端的常量内存数组（bilateral_kernel.cu），上限 64 个元素。由于 2D 高斯核在行、列方向对称，示例只生成一维高斯掩码，再通过行列索引乘积得到二维权重（见 updateGaussian）：

extern "C" void updateGaussian(float delta, int radius) { float fGaussian[64]; for (int i = 0; i < 2 * radius + 1; ++i) { float x = (float)(i - radius); fGaussian[i] = expf(-(x * x) / (2 * delta * delta)); } checkCudaErrors(cudaMemcpyToSymbol(cGaussian, fGaussian, sizeof(float) * (2 * radius + 1))); }

实现要点：源码注释明确指出，将一维高斯数组映射为 1D 纹理而非共享内存，是为了避免共享内存访问可能引发的严重 bank conflict；同时线程采用 y 方向（列方向）遍历，保证全局内存写入是合并（coalesced）的（见 bilateral_kernel.cu）。

3.2 纹理对象：缓存命中的关键

示例通过cudaCreateTextureObject创建两个 2D 纹理对象（initTexture）：

• rgbaTexdImage：绑定原始图像，filterMode = cudaFilterModePoint（点采样，不做插值）；
• rgbaTexdTemp：绑定迭代用的临时缓冲区，filterMode = cudaFilterModeLinear（线性过滤）。

两者均使用cudaResourceTypePitch2D类型的资源描述，通道格式为uchar4，读取模式为cudaReadModeNormalizedFloat，即读取时自动将 8 位整数归一化到[0, 1]浮点区间。由于是逐像素非线性处理，源码注释明确说明该示例"只是简单的逐像素步进"，未做共享内存分块优化；纹理硬件自带的边界钳制（clamp to edge）行为则省去了手动处理越界索引的工作。

3.3 迭代执行与多趟滤波

bilateralFilterRGBA 是宿主端调度函数，它按iterations次数循环启动 kernel：

• 迭代次数为 1 时，直接从rgbaTexdImage读取；
• 迭代次数大于 1 时，从rgbaTexdTemp读取（前一轮结果），每轮结束后通过cudaMemcpy2D（DeviceToDevice）把结果拷回dTemp，作为下一轮输入，从而形成迭代式滤波。

kernel 的启动配置为固定16 × 16线程块，网格按图像尺寸向上取整（bilateral_kernel.cu）。同时该函数会统计每轮 kernel 的执行耗时（通过StopWatchInterface计时器），并返回平均 kernel 时间，供基准测试模式使用。

3.4 CPU 黄金参考实现（验证用）

bilateralFilter_cpu.cpp 提供了 CPU 端的参考实现bilateralFilterGold与updateGaussianGold，用于验证 GPU 结果的正确性。它逐像素遍历(2r+1)²邻域，对越界的邻居坐标做显式钳制（clamp），采用与 GPU kernel 相同的权重公式。GPU 输出会与 CPU 参考结果通过sdkComparePPM比对（容差由MAX_EPSILON_ERROR 5.0f和 0.15 的阈值控制，见 bilateralFilter.cpp）。

四、CUDA-OpenGL 图形互操作管线

4.1 数据流：GPU 滤波 → PBO → 纹理显示

该示例同时是"Graphics Interop"关键概念的典型示范。整个显示管线（display 函数）如下：

1. 调用cudaGraphicsMapResources映射 OpenGL 像素缓冲对象（PBO）；
2. 通过cudaGraphicsResourceGetMappedPointer拿到 PBO 对应的设备指针dResult；
3. 执行bilateralFilterRGBA把滤波结果直接写入 PBO；
4. 调用cudaGraphicsUnmapResources解除映射，交还 OpenGL；
5. OpenGL 端把 PBO 作为GL_PIXEL_UNPACK_BUFFER_ARB绑定，通过glTexSubImage2D上传为 2D 纹理；
6. 绑定 ARB 片段程序（汇编着色器!!ARBfp1.0，仅一行TEX指令，见 bilateralFilter.cpp）后以全屏四边形绘制。

PBO 的注册发生在 initGLResources：cudaGraphicsGLRegisterBuffer(&cuda_pbo_resource, pbo, cudaGraphicsMapFlagsWriteDiscard)，其中WriteDiscard标志告诉 CUDA 将独占写入该缓冲，从而获得更好的性能。程序退出时通过cudaGraphicsUnregisterResource注销资源。

4.2 渲染环境要求

initGL 中会校验 OpenGL 版本与扩展：

• OpenGL 2.0 及以上；
• GL_ARB_vertex_buffer_object、GL_ARB_pixel_buffer_object扩展；
• GLUT 窗口初始化使用GLUT_RGBA | GLUT_DOUBLE显示模式，窗口尺寸与输入图像一致。

五、图像数据与加载

示例默认加载nature_monte.bmp（位于 cpp/5_Domain_Specific/bilateralFilter/data/nature_monte.bmp，640×480）。这是一张包含苹果、海螺壳与秋叶的静物照片，画面同时具备光滑表面（苹果）、细腻纹理（海螺壳螺旋纹、叶脉）与粗糙背景（石质表面），非常适合直观演示双边滤波"平滑纹理噪声、保留边缘轮廓"的特性。

图像加载由自实现的 bmploader.cpp 完成。LoadBMPFile支持标准的 24 位（bitsPerPixel == 24）、非压缩 BMP 文件：读取 BMP 头与信息头后，按行逐像素读取 BGR 三通道并补齐每行的 4 字节对齐填充，转换为uchar4（w 分量置 0）输出。由于不使用系统库，该加载器对 BMP 格式做了严格校验（文件魔数0x4D42、颜色深度、压缩标志），不满足条件即报错退出。

数据目录中还包含若干参考输出图（ref_05.ppm、ref_06.ppm、ref_07.ppm、ref_08.ppm），分别对应不同滤波半径下的期望结果，供自动验证模式比对使用。

六、构建、运行与三种运行模式

6.1 前置依赖与支持范围

根据示例 README 与 CMakeLists.txt，构建/运行需要：

• CUDA Toolkit：请先为对应平台下载安装 CUDA Toolkit；
• X11（Linux 图形环境）；
• OpenGL（2.0+，含 PBO/VBO 扩展）；
• FreeGLUT与GLEW：Windows 下直接使用仓库 Common/lib/x64 中的freeglut.lib、glew64.lib（若存在glew32.lib则自动改用 glew32），Linux 下通过find_package(OpenGL)/find_package(GLUT)探测。

官方 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。

CMake 构建脚本默认将CMAKE_CUDA_ARCHITECTURES设为75 80 86 87 89 90 100 110 120，并要求 C++/CUDA 标准为 C++17；若未找到 OpenGL 或 GLUT，构建脚本会打印提示并跳过该示例（见 CMakeLists.txt）。

6.2 命令行参数

在 main 中，示例支持以下参数（见 printHelp）：

6.3 三种运行模式

模式一：默认 OpenGL 交互模式（不带-radius/-passes/-benchmark） 初始化 GLUT 窗口与 CUDA 资源后进入glutMainLoop，以 10ms 间隔（REFRESH_DELAY）刷新画面，窗口标题实时显示 FPS 与当前参数。非交互（动画开）状态下，varyEuclidean会以 1.02 倍因子自动往复调节 euclidean delta（bilateralFilter.cpp），演示参数变化对画面的影响。

模式二：单次测试模式（指定-radius或-passes） 调用runSingleTest：执行一次滤波、将结果写入nature_%02d.ppm，再与-file指定的参考图用sdkComparePPM比对，输出 "Image is Matching / Different"，并以此决定进程退出码。

模式三：基准测试模式（-benchmark） 调用runBenchmark：先做一次 warm-up，再循环执行 150 次滤波，统计平均 kernel 时间并输出吞吐量：

bilateralFilter-texture, Throughput = X.XXXX M RGBA Pixels/s, Time = X.XXXXX s, Size = XXXXX RGBA Pixels, NumDevsUsed = 1

七、涉及的核心 CUDA Runtime API

示例 README 列出的 CUDA Runtime API 及其在本示例中的用途可汇总如下：

八、小结

bilateralFilter 示例以一份紧凑的代码展示了三层核心能力：

1. 算法层：完整实现 Tomasi & Manduchi 双边滤波——空间域高斯核 × 颜色域欧氏权重的加权平均，并通过三个交互参数直观演示"去纹理、保边缘、迭代出卡通效果"的差异化表现；
2. 性能层：利用常量内存、纹理对象（点采样与线性过滤双纹理）、合并的列方向遍历与 PBO 直写，构建了低开销的 GPU 滤波管线；
3. 工程层：提供 CPU 黄金参考实现与 PPM 自动比对、150 轮基准测试、命令行参数控制三种运行模式，兼顾教学验证与性能测量。

对希望在 CUDA 中实现图像恢复、实时去噪或"边缘保持平滑"类后处理效果的开发者而言，bilateral_kernel.cu 的 kernel 与纹理用法、bilateralFilter.cpp 的 PBO 互操作管线，都是可直接借鉴的范本；而同目录下的 ref_05.ppm 等参考输出则可用于自行扩展验证。

【免费下载链接】cuda-samplesSamples for CUDA Developers which demonstrates features in CUDA Toolkit项目地址: https://gitcode.com/GitHub_Trending/cu/cuda-samples