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 SADSum 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*RAD1) x (2*RAD1) 17 x 17。对于像素(x, y)和视差d代价定义为cost(x, y, d) Σ Σ |left(xj, yi) - right(xjd, yi)| (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 之和。从源码注释可以确认该指令自 KeplerSM 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-L124for (int i 0; i STEPS; i) { int offset -RAD i * RAD; imLeftA[i] tex2Dunsigned int(tex2Dleft, tidx - RAD, tidy offset); imLeftB[i] tex2Dunsigned int(tex2Dleft, tidx - RAD blockSize_x, tidy offset); }imLeftA服务于本线程的列imLeftB服务于右邻blockSize_x列供相邻线程协作完成右边界像素的窗口计算。左图数据只读一次、复用于全部 17 个视差候选大幅削减了纹理内存带宽消耗。视差循环内的三阶段流水对每个视差dminDisparity到maxDisparity示例共 17 个候选kernel 依次执行计算 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采样。水平窗口求和在共享内存中沿x方向对[-RAD, RAD]区间累加并在cg::sync(cta)栅栏同步下原地回写stereoDisparity_kernel.cuh#L156-L171。垂直窗口求和与最优选择沿y方向对[-RAD, RAD]区间累加得到最终代价cost与历史最优bestCost比较并更新bestDisparity d 8stereoDisparity_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。使用sdkFindFilePathsdkLoadPPM4ub加载数据目录下的左右视图 stereo.im0.640x533.ppm 与 stereo.im1.640x533.ppm图像按 4 字节/像素RGBX读入尺寸为 640x533stereoDisparity.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描述符为cudaCreateChannelDescunsigned 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*RAD1得出、视差搜索范围[-16:0]GPU 处理时间毫秒与像素吞吐率Mpixels/sec(w * h * 1000) / msecTotal / 1000000。结果输出与 CPU 校验计算完成后将d_odata拷回 Host计算 GPU 输出视差图所有像素的Checksum将视差值乘以缩放因子mult 20后以 PGM 格式写出output_GPU.pgmstereoDisparity.cu#L234-L244调用 CPU 参考实现cpu_gold_stereo重新计算一遍得到 CPU Checksum 并写出output_CPU.pgm程序以exit((checkSum cpuCheckSum) ? EXIT_SUCCESS : EXIT_FAILURE)结束stereoDisparity.cu#L284即GPU 结果与 CPU 参考解必须逐位一致否则视为失败退出。CPU 参考实现cpu_gold_stereostereoDisparity_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时切换为-Gcuda-gdb 调试通过target_compile_features启用 C17 与 CUDA 17 标准并开启CUDA_SEPARABLE_COMPILATIONPOST_BUILD 阶段自动把data目录拷贝到构建目录保证运行时能定位输入 PPM 图像最后引入InstallSamples.cmake以支持安装到 CUDA Samples 标准目录。典型构建命令与仓库其他示例一致mkdir build cd build cmake .. -DCMAKE_BUILD_TYPERelease make -j ./stereoDisparity运行输出与产物程序运行时依次打印设备信息、加载的输入文件、kernel 配置、GPU 处理时间与像素吞吐率、GPU/CPU Checksum 以及输出文件路径。运行结束后当前目录会生成output_GPU.pgmGPU 计算的视差图视差值 × 20 缩放后保存output_CPU.pgmCPU 参考实现的视差图。由于程序内建了 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创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考