企业官网建设流程全解析

CUDA Samples 视差计算实战：基于 SIMD SAD 视频内建指令的 stereoDisparity 立体匹配示例

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

导读

本文深入剖析 CUDA Samples 仓库（cpp/5_Domain_Specific 分类）中的 stereoDisparity 示例，它演示了如何使用 CUDA 的 SIMD SAD（Sum of Absolute Difference，绝对差之和）视频内建指令在 GPU 上并行计算立体视差图（stereo disparity map）。通过阅读本文，你将理解块匹配（block matching）视差算法的原理、vabsdiff4视频 intrinsics 的 PTX 汇编细节、共享内存与寄存器预取优化手段，以及示例自带的 CPU 参考实现与校验流程，并掌握该示例的完整构建、运行与结果验证方法。

示例概览与关键概念

stereoDisparity 是一个独立的 CUDA 程序，其核心目标是对一对左右立体图像（left/right image pair）计算视差图。视差（disparity）反映了同一场景点在左右两幅图像中成像位置的水平偏移量，是双目立体视觉中恢复深度信息的关键中间量：偏移越大意味着视差越大，物体距离相机越近。

• 关键概念：Image Processing（图像处理）、Video Intrinsics（视频内建指令）。
• 支持的操作系统：Linux、Windows。
• 支持的 CPU 架构：x86_64、armv7l。
• 支持的 SM 架构：SM 5.0、SM 5.2、SM 5.3、SM 6.0、SM 6.1、SM 7.0、SM 7.2、SM 7.5、SM 8.0、SM 8.6、SM 8.7、SM 8.9、SM 9.0。
• 前置条件：在对应平台下载并安装 CUDA Toolkit。

示例使用的 CUDA Runtime API 包括：cudaMemcpy、cudaFree、cudaEventSynchronize、cudaDeviceSynchronize、cudaCreateTextureObject、cudaEventRecord、cudaMalloc、cudaEventElapsedTime、cudaGetDeviceProperties、cudaEventCreate。

算法原理：块匹配与 SAD 代价

视差搜索的基本思想

对于左图中的每一个像素，算法在其水平搜索范围内（示例中为minDisp = -16到maxDisp = 0）尝试不同的水平位移d，将左图以该像素为中心的一个窗口与右图中水平偏移d后的对应窗口进行比较，计算两者之间的匹配代价。代价最小的位移d就被认定为该像素的视差值。

在 stereoDisparity_kernel.cuh 中，算法的表述非常清晰：

对左图中候选像素的一个区域与右图的不同水平移位区域之间计算绝对差之和；取差值和最小的那个位移作为该像素在左右图像对之间的移动量。恢复出的运动量就是视差图。运动越多意味着视差越大，即物体越近。

SAD 代价的数学形式

窗口大小由半径RAD决定，示例中RAD = 8，因此匹配窗口为(2*RAD+1) x (2*RAD+1) = 17 x 17。对于像素(x, y)和视差d，代价定义为：

cost(x, y, d) = Σ Σ |left(x+j, y+i) - right(x+j+d, y+i)| (i, j ∈ [-RAD, RAD])

SIMD 视频内建指令__usad4

示例的关键亮点在于用 SIMD 指令一次完成四个 8-bit 通道的绝对差求和。输入图像以 4 字节/像素（RGBA/RGBX）存储，即一个unsigned int中打包了 4 个 8-bit 分量。__usad4用内联 PTX 汇编实现（见 stereoDisparity_kernel.cuh#L57-L68）：

__device__ unsigned int __usad4(unsigned int A, unsigned int B, unsigned int C = 0) { unsigned int result; // Kepler (SM 3.x) and higher supports a 4 vector SAD SIMD asm("vabsdiff4.u32.u32.u32.add" " %0, %1, %2, %3;" : "=r"(result) : "r"(A), "r"(B), "r"(C)); return result; }

vabsdiff4.u32.u32.u32.add是 PTX 指令集中的 4 路向量绝对差求和指令：它将A与B按 4 个 8-bit 通道分别计算绝对差，再把结果按 8-bit 通道对齐累加（由.add后缀完成通道内加法），最终与累加器C相加，得到四通道 SAD 之和。从源码注释可以确认，该指令自 Kepler（SM 3.x）及更高架构起可用；对应文档可参考 CUDA Toolkit 附带的《Using_Inline_PTX_Assembly_In_CUDA》与相应架构的《PTX ISA》手册。

相比标量逐个字节求差再累加，__usad4将 4 个通道的计算压缩为一条 SIMD 指令，显著提升了匹配代价计算的数据吞吐，这正是该示例被称为“SAD SIMD Intrinsics”的原因。

核心 Kernel 实现剖析

KernelstereoDisparityKernel定义在 stereoDisparity_kernel.cuh#L90-L193，是整篇算法的 GPU 端主体。其设计包含多个值得学习的优化点。

线程块与共享内存布局

#define blockSize_x 32 #define blockSize_y 8 #define RAD 8 // 搜索支持区域的半径 #define STEPS 3 // 初始化共享内存所需的加载次数

线程块大小为32 x 8。由于每个像素的 17 x 17 匹配窗口会跨越相邻线程，kernel 在共享内存中维护一个带RAD边界扩展的代价缓存：

__shared__ unsigned int diff[blockSize_y + 2 * RAD][blockSize_x + 2 * RAD];

即48 x 48的共享内存数组，用于暂存每个视差d下各像素的 SAD 代价，随后在水平与垂直两个方向上进行窗口累加。

寄存器预取左图像素

为减少纹理访问次数，kernel 先将左图需要的数据预取到寄存器数组imLeftA[STEPS]与imLeftB[STEPS]中（STEPS = 3，对应-RAD、0、RAD三个纵向偏移，见 stereoDisparity_kernel.cuh#L120-L124）：

for (int i = 0; i < STEPS; i++) { int offset = -RAD + i * RAD; imLeftA[i] = tex2D<unsigned int>(tex2Dleft, tidx - RAD, tidy + offset); imLeftB[i] = tex2D<unsigned int>(tex2Dleft, tidx - RAD + blockSize_x, tidy + offset); }

imLeftA服务于本线程的列，imLeftB服务于右邻blockSize_x列（供相邻线程协作完成右边界像素的窗口计算）。左图数据只读一次、复用于全部 17 个视差候选，大幅削减了纹理内存带宽消耗。

视差循环内的三阶段流水

对每个视差d（minDisparity到maxDisparity，示例共 17 个候选），kernel 依次执行：

1. 计算 SAD 并写入共享内存：LEFT 阶段，用__usad4(imLeft, imRight)计算本线程列的代价；RIGHT 阶段，threadIdx.x < 2*RAD的线程额外计算右边界blockSize_x处的代价（stereoDisparity_kernel.cuh#L127-L152）。右图通过tex2D(tex2Dright, tidx - RAD + d, tidy + offset)以水平偏移d采样。
2. 水平窗口求和：在共享内存中沿x方向对[-RAD, RAD]区间累加，并在cg::sync(cta)栅栏同步下原地回写（stereoDisparity_kernel.cuh#L156-L171）。
3. 垂直窗口求和与最优选择：沿y方向对[-RAD, RAD]区间累加得到最终代价cost，与历史最优bestCost比较并更新bestDisparity = d + 8（stereoDisparity_kernel.cuh#L173-L188）。偏移+8用于把视差范围映射到非负值。

整个内核使用 cooperative_groups 的cg::this_thread_block()与cg::sync(cta)进行块内同步，代码中对所有内层循环都标注了#pragma unroll以利于编译器展开。

最后，在边界检查tidy < h && tidx < w通过后，将bestDisparity写入全局输出g_odata[tidy * w + tidx]。

主程序：数据加载、纹理对象与性能计时

主程序 stereoDisparity.cu 负责完整的“加载 → 计算 → 计时 → 验证 → 输出”流程。

图像加载与设备内存

• 通过findCudaDevice自动选择最佳 CUDA 设备，并打印 SM 版本与多处理器数量（stereoDisparity.cu#L76-L84）。
• 使用sdkFindFilePath+sdkLoadPPM4ub加载数据目录下的左右视图 stereo.im0.640x533.ppm 与 stereo.im1.640x533.ppm，图像按 4 字节/像素（RGBX）读入，尺寸为 640x533（stereoDisparity.cu#L97-L113）。
• 网格规模按iDivUp向上取整计算：numThreads = (32, 8, 1)，numBlocks = (iDivUp(w,32), iDivUp(h,8))（stereoDisparity.cu#L115-L116）。
• 分别cudaMalloc输出视差图d_odata与两幅输入图d_img0、d_img1，并用cudaMemcpy完成 Host 到 Device 的拷贝（stereoDisparity.cu#L127-L137）。

纹理对象（Texture Object）配置

kernel 通过纹理对象而非全局内存直接访问输入图像，以利用纹理缓存的局部性与硬件采样支持。配置要点（stereoDisparity.cu#L139-L181）：

• 资源类型为cudaResourceTypePitch2D，绑定到设备指针d_img0/d_img1，描述符为cudaCreateChannelDesc<unsigned int>()，pitch 为w * 4字节；
• 纹理描述：normalizedCoords = false（非归一化坐标）、filterMode = cudaFilterModePoint（最近邻采样，保证不做插值）、addressMode[0/1] = cudaAddressModeClamp（越界钳制到边界像素）、readMode = cudaReadModeElementType；
• 左右两幅图像分别通过cudaCreateTextureObject创建tex2Dleft与tex2Dright。

其中cudaAddressModeClamp相当于在 GPU 端天然实现了 CPU 参考代码里的边界钳制逻辑，避免了显式的越界分支。

预热运行与 CUDA Event 计时

为保证测量稳定，代码先执行一次“预热 kernel”（warmup run），使 GPU 进入最大功耗状态（stereoDisparity.cu#L183-L187）。随后创建cudaEvent_t事件对，通过cudaEventRecord(start)→ 启动 kernel →cudaEventRecord(stop)→cudaEventSynchronize(stop)→cudaEventElapsedTime得到 GPU 处理耗时（stereoDisparity.cu#L189-L213）。程序会打印：

• 输入尺寸（640x533）、kernel 窗口尺寸（17x17，由2*RAD+1得出）、视差搜索范围[-16:0]；
• GPU 处理时间（毫秒）与像素吞吐率（Mpixels/sec）：(w * h * 1000) / msecTotal / 1000000。

结果输出与 CPU 校验

计算完成后：

1. 将d_odata拷回 Host，计算 GPU 输出视差图所有像素的Checksum；
2. 将视差值乘以缩放因子mult = 20后以 PGM 格式写出output_GPU.pgm（stereoDisparity.cu#L234-L244）；
3. 调用 CPU 参考实现cpu_gold_stereo重新计算一遍，得到 CPU Checksum 并写出output_CPU.pgm；
4. 程序以exit((checkSum == cpuCheckSum) ? EXIT_SUCCESS : EXIT_FAILURE)结束（stereoDisparity.cu#L284），即GPU 结果与 CPU 参考解必须逐位一致，否则视为失败退出。

CPU 参考实现

cpu_gold_stereo（stereoDisparity_kernel.cuh#L195-L263）是 GPU kernel 的朴素串行版本：对每个像素、每个视差、窗口内每个位置三重循环，逐字节求绝对差（abs((int)(A[k] - B[k]))对 4 个通道求和），并对越界坐标显式钳制到[0, w-1]/[0, h-1]。它的存在既作为正确性参照物，也直观展示了 SIMD 指令相对朴素实现的性能优势来源。

构建与运行

构建方式

示例提供独立的 CMakeLists.txt，要点如下：

• 要求 CMake 3.20+，语言集合为C CXX CUDA，通过find_package(CUDAToolkit REQUIRED)定位 CUDA；
• 默认目标架构列表为75 80 86 87 89 90 100 110 120，可通过CMAKE_CUDA_ARCHITECTURES覆盖；
• 编译选项默认携带-lineinfo（加入行号信息便于调试工具），开启ENABLE_CUDA_DEBUG时切换为-G（cuda-gdb 调试）；
• 通过target_compile_features启用 C++17 与 CUDA 17 标准，并开启CUDA_SEPARABLE_COMPILATION；
• POST_BUILD 阶段自动把data目录拷贝到构建目录，保证运行时能定位输入 PPM 图像；
• 最后引入InstallSamples.cmake以支持安装到 CUDA Samples 标准目录。

典型构建命令（与仓库其他示例一致）：

mkdir build && cd build cmake .. -DCMAKE_BUILD_TYPE=Release make -j ./stereoDisparity

运行输出与产物

程序运行时依次打印设备信息、加载的输入文件、kernel 配置、GPU 处理时间与像素吞吐率、GPU/CPU Checksum 以及输出文件路径。运行结束后，当前目录会生成：

• output_GPU.pgm：GPU 计算的视差图（视差值 × 20 缩放后保存）；
• output_CPU.pgm：CPU 参考实现的视差图。

由于程序内建了 Checksum 一致性校验，只要最终退出码为 0，即可确认 GPU kernel 与 CPU 参考实现的计算结果完全一致，这也为将该示例作为其它视差算法移植的起点提供了可靠的回归验证手段。

总结

stereoDisparity 是一个“小而精”的 CUDA 示例：它以 17 x 17 块匹配 + 17 级视差搜索为算法骨架，用vabsdiff4SIMD 视频内建指令压缩四通道 SAD 计算，用共享内存 + 寄存器预取 + cooperative groups 栅栏同步组织并行数据流，并用纹理对象的 Clamp 寻址天然处理边界，最后以 GPU/CPU 双 Checksum 校验保障正确性。无论你是学习立体视觉算法、视频 intrinsics 内联汇编，还是研究共享内存优化的经典模式，这份示例源码都是值得逐行研读的范本。

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