CUDA 盒子滤波（Box Filter）示例深度解析：滑动窗口算法与 OpenGL 互操作实战

CUDA 盒子滤波（Box Filter）示例深度解析：滑动窗口算法与 OpenGL 互操作实战

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

本篇文章以 CUDA Samples 仓库中的 boxFilter 示例为主体，系统讲解如何在 CUDA 上实现一个与滤波半径无关的常量开销快速盒子滤波，并借助 OpenGL 像素缓冲对象（PBO）实时渲染结果。读完本文，你将掌握滑动窗口累加算法、可分离滤波的两遍（pass）实现、CUDA 纹理对象与全局内存的访问模式权衡，以及 CUDA–OpenGL 图形互操作的标准流程，并能在自己的项目中复刻这套"滤波 + 可视化 + 自动验证"的完整范例。

示例概览：一个集滤波、渲染、验证于一体的 CUDA 图像处理样例

boxFilter 是 CUDA Samples 仓库中2_Concepts_and_Techniques目录下的图像处理示例（仓库内其余概念与技术示例见 cpp/2_Concepts_and_Techniques/README.md），其 README 给出的定位是：

Fast image box filter using CUDA with OpenGL rendering.

即：用 CUDA 对图像做快速盒子滤波（box filter），并用 OpenGL 实时显示结果。该示例覆盖两大关键技术主题（Key Concepts）：Graphics Interop（图形互操作）与Image Processing（图像处理）。

从源码头注释（boxFilter.cpp）可以进一步看到它的几个核心设计点：

• 行（rows）与列（columns）被并行处理；
• 盒子滤波被实现为常量开销（constant cost）算法，与滤波宽度无关；
• 默认处理 8 位 RGBA 图像（Version 1.1 起）；
• 通过键盘=/-可动态增减滤波半径，直观观察模糊效果变化。

示例运行时加载的默认图像为teapot1024.ppm（数据文件位于 cpp/2_Concepts_and_Techniques/boxFilter/data），默认滤波半径filter_radius = 14、迭代次数iterations = 1、线程数nthreads = 64。

算法原理：滑动窗口（Sliding Window）实现常量开销滤波

传统盒子滤波需要对窗口内每个像素求和，复杂度随半径线性增长。而 boxFilter 采用滑动窗口法：当滤波核从左向右滑动时，只需"加入右侧新进入窗口的采样值，减去左侧滑出窗口的采样值"。

boxFilter_kernel.cu 的注释精确描述了这一特性：

As the kernel moves from left to right, we add in the contribution of the new sample on the right, and subtract the value of the exiting sample on the left. This only requires 2 adds and a mul per output value, independent of the filter radius.

即每个输出像素仅需 2 次加法 + 1 次乘法，与滤波半径 r 完全无关。以d_boxfilter_x（boxFilter_kernel.cu）为例，其核心循环为：

float scale = 1.0f / (float)((r << 1) + 1); float t; // 左边缘初始化：复制边缘像素，保证窗口越界时依然有值 t = id[0] * r; for (int x = 0; x < (r + 1); x++) { t += id[x]; } od[0] = t * scale; // 主循环：滑动窗口，只做加一减一 for (int x = (r + 1); x < w - r; x++) { t += id[x + r]; // 加入窗口右侧新像素 t -= id[x - r - 1]; // 减去窗口左侧滑出像素 od[x] = t * scale; }

窗口宽度为(r << 1) + 1，即2r+1，归一化系数scale = 1/(2r+1)。边缘处理采用复制边缘像素策略（"This version duplicates edge pixels"），在左右（上下）边界通过重复首尾像素避免越界。

可分离滤波：水平 + 垂直两遍扫描

盒子滤波是**可分离（separable）**的：2D 盒子滤波可以分解为先做水平方向（x pass）、再做垂直方向（y pass）两个一维滤波。两遍之间的中间结果存放在设备端临时缓冲d_temp中。

一维滤波分别在d_boxfilter_x与d_boxfilter_y（boxFilter_kernel.cu）中实现，结构完全对称。以 8 位 RGBA 数据路径为例，主机端驱动函数boxFilterRGBA（boxFilter_kernel.cu）按如下顺序执行：

1. 水平 pass：d_boxfilter_rgba_x<<<height / nthreads, nthreads, 0>>>(...)，使用纹理查找；
2. 垂直 pass：d_boxfilter_rgba_y<<<width / nthreads, nthreads, 0>>>(...)，使用全局内存；
3. 若迭代次数iterations > 1，通过cudaMemcpy2DToArray（cudaMemcpyDeviceToDevice）把结果拷回纹理临时数组d_tempArray，供下一轮迭代的水平 pass 使用。

多轮迭代的意义在源码注释中也有明确说明：对盒子滤波多次迭代会逐渐收敛于高斯模糊（"Applying multiple iterations of the box filter converges towards a Gaussian blur"），这也是该示例提供passes参数的原因——用若干次廉价盒子滤波近似高斯效果。

数据路径：浮点灰度版与 8 位 RGBA 版

boxFilter 同时提供两条数据路径：

RGBA 路径把 32 位无符号整数视为 4 个 8 位通道（R/G/B/A），并在内核中通过rgbaIntToFloat/rgbaFloatToInt（boxFilter_kernel.cu）完成与浮点float4的转换，转换时使用__saturatef将颜色钳制到[0.0, 1.0]区间，再按×255回写 8 位通道。默认交互模式即走 RGBA 路径（initCuda(true)）。

内存访问优化：纹理查找 vs 全局内存合并

d_boxfilter_x的注释揭示了一个重要的 CUDA 性能细节：

Note that the x (row) pass suffers from uncoalesced global memory reads, since each thread is reading from a different row. For this reason it is better to use texture lookups for the x pass. The y (column) pass is perfectly coalesced.

水平 pass 中每个线程读取一行、行与行之间内存不连续，全局内存访问无法合并，因此 x pass 采用纹理对象查找（tex2D<float4>/tex2D<float>）；垂直 pass 中每列线程按id[y * w]步进访问、同一 warp 内相邻线程访问连续地址，天然合并，因而直接使用全局内存。这是"用纹理缓存放宽未合并访问、把合并访问留给全局内存"的典型优化组合。

纹理对象的构建集中在initTexture（boxFilter_kernel.cu），通过cudaResourceDesc+cudaTextureDesc+cudaCreateTextureObject创建：

• 输入纹理rgbaTex：normalizedCoords = false、filterMode = cudaFilterModeLinear、地址模式cudaAddressModeWrap、readMode = cudaReadModeNormalizedFloat；
• 临时缓冲纹理rgbaTexTempArray：地址模式改为cudaAddressModeClamp（配合边缘复制策略，越界采样自动钳制到边缘像素）；
• 浮点路径的tex/texTempArray则使用cudaFilterModePoint、cudaReadModeElementType。

这里cudaCreateTextureObject对应 README 中列出的cudaCreateChannelDesc、cudaMallocArray、cudaCreateTextureObject、cudaDestroyTextureObject、cudaMemcpy2DToArray（经cudaMemcpy族）等一系列 CUDA Runtime API。

图形互操作：PBO + CUDA 的完整渲染管线

README 的核心概念之一是Graphics Interop。示例在 boxFilter.cpp 中实现了标准的三步式 CUDA–OpenGL 互操作流程：

1. 注册：cudaGraphicsGLRegisterBuffer(&cuda_pbo_resource, pbo, cudaGraphicsMapFlagsWriteDiscard)（boxFilter.cpp），将 OpenGL 像素缓冲对象 PBO 注册为 CUDA 图形资源；
2. 映射：每帧渲染前cudaGraphicsMapResources+cudaGraphicsResourceGetMappedPointer拿到设备端指针d_result，直接作为滤波输出目标（boxFilter.cpp）；
3. 解除映射：cudaGraphicsUnmapResources后交由 OpenGL 使用——glTexSubImage2D从 PBO 上传纹理，再用 ARB 汇编片元着色器!!ARBfp1.0（boxFilter.cpp）绘制浮点纹理并交换缓冲（glutSwapBuffers）。

渲染主循环由glutMainLoop驱动，timerEvent每REFRESH_DELAY = 10ms 触发一次重绘；程序退出时cleanup（boxFilter.cpp）统一释放纹理对象（freeTextures→cudaDestroyTextureObject/cudaFreeArray）、设备缓冲（cudaFree）并cudaGraphicsUnregisterResource注销 PBO。这与 README 中列出的 API 清单（cudaGraphicsMapResources、cudaGraphicsResourceGetMappedPointer、cudaGraphicsUnmapResources、cudaGraphicsUnregisterResource、cudaGraphicsGLRegisterBuffer等）一一对应。

启动时initGL（boxFilter.cpp）还会校验 OpenGL 版本与扩展：要求OpenGL 2.0、GL_ARB_vertex_buffer_object、GL_ARB_pixel_buffer_object，不满足则直接退出。

交互控制与自动动画

运行默认模式后，窗口标题实时显示CUDA Rolling Box Filter <Animation=ON> (radius=…, passes=…): … fps。键盘控制（keyboard回调，boxFilter.cpp）汇总如下：

在自动动画模式下，varySigma（boxFilter.cpp）会让半径在 0～64 之间往复摆动，配合computeFPS（boxFilter.cpp）的 FPS 统计，可直观看到半径变化对模糊强度与吞吐的影响。

命令行参数与基准测试

main（boxFilter.cpp）支持以下参数（-help可查看printHelp输出）：

三种运行模式由参数决定：

1. 交互可视化模式（默认）：启动 GLUT 窗口，实时渲染并支持键盘调节；
2. 单测校验模式：指定-radius或-passes时进入runSingleTest（boxFilter.cpp）——滤波结果写为teapot1024_%02d.ppm，再与参考图ref_14.ppm/ref_22.ppm比对；
3. 基准测试模式：-benchmark进入runBenchmark（boxFilter.cpp）——先做一次 warm-up，再连续执行 150 轮（iCycles = 150）RGBA 滤波，输出吞吐率Throughput = … M RGBA Pixels/s与平均耗时。

校验环节使用sdkComparePPM，容差为MAX_EPSILON_ERROR = 5.0f、相对误差0.15f（boxFilter.cpp），参考图即数据目录中的 ref_14.ppm 与 ref_22.ppm。而 CPU 黄金参考实现computeGold（水平 + 垂直两遍，逻辑与 GPU 内核一一对应）位于 boxFilter_cpu.cpp，可用作正确性的交叉验证。

构建与运行

该示例的构建目标定义在 CMakeLists.txt，要点如下：

• 源文件：boxFilter.cpp、boxFilter_cpu.cpp、boxFilter_kernel.cu，启用cxx_std_17/cuda_std_17与CUDA_SEPARABLE_COMPILATION；
• 默认编译目标架构：CMAKE_CUDA_ARCHITECTURES为75 80 86 87 89 90 100 110 120（对应 SM 7.5 至 SM 12.0）；
• 依赖find_package(OpenGL)与find_package(GLUT)，两者任一缺失时该示例不会编译（仅打印提示）；
• Windows 下额外链接Common/lib/x64中的freeglut.lib与glew64.lib（ARM 上为glew32.lib），并将 DLL 拷贝到输出目录；
• 构建后自动把data目录复制到输出目录，保证运行时能通过sdkFindFilePath找到teapot1024.ppm。

按仓库根 README.md 的通用流程，可在克隆仓库后执行：

mkdir build && cd build cmake .. make -j$(nproc)

然后到build/cpp/2_Concepts_and_Techniques/boxFilter目录运行./boxFilter即可进入可视化模式；或运行./boxFilter -radius=14 -passes=1做单测校验、./boxFilter -benchmark跑性能基准。Linux 下示例通过 X11 建立显示环境（源码中setenv("DISPLAY", ":0", 0)），因此需要可用的图形会话；依赖项（X11、OpenGL、Freeglut、GLEW）的说明与安装指引见仓库根 README.md 的 Dependencies 一节。

支持矩阵与 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 Runtime API 完整清单（摘自 README）为：cudaGraphicsUnmapResources、cudaCreateChannelDesc、cudaMallocArray、cudaFreeArray、cudaFree、cudaGetErrorString、cudaMemcpy、cudaGraphicsResourceGetMappedPointer、cudaGraphicsMapResources、cudaDestroyTextureObject、cudaDeviceSynchronize、cudaCreateTextureObject、cudaGraphicsUnregisterResource、cudaMalloc、cudaGraphicsGLRegisterBuffer——它们分别服务于纹理对象构建、设备内存管理、图形资源映射与错误处理四大环节，构成了一条完整的"纹理滤波 + 图形互操作"代码路径。

小结

boxFilter 示例麻雀虽小五脏俱全：它以滑动窗口 + 可分离滤波实现了与半径无关的 O(1) 盒子滤波，用"纹理处理水平 pass、全局内存处理垂直 pass"的混合访问策略兼顾了访存效率，再通过 PBO 图形互操作把 CUDA 计算结果实时呈现到 OpenGL 窗口，并配套 CPU 黄金参考、参考图比对与 150 轮基准测试三套验证手段。无论是想学习 CUDA 图像滤波优化、纹理对象 API，还是 CUDA–OpenGL 互操作的工程范式，这个样例都是理想的参考蓝本。

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