基于纹理对象（Texture Object）的高斯可分离卷积：CUDA Samples 的 convolutionTexture 深入解析

基于纹理对象Texture Object的高斯可分离卷积CUDA Samples 的 convolutionTexture 深入解析【免费下载链接】cuda-samplesSamples for CUDA Developers which demonstrates features in CUDA Toolkit项目地址: https://gitcode.com/GitHub_Trending/cu/cuda-samplesCUDA 样本 convolutionTexture 演示了如何完全使用纹理Texture而非共享内存实现二维高斯核的可分离卷积其定位是与同目录下的 convolutionSeparable 样本进行性能对比。读完本文你将掌握纹理对象的创建与绑定、tex2D纹理采样在卷积核中的用法、模板展开loop unrolling优化技巧以及一套完整的 GPU 结果对 CPU 参考实现做 L2 范数校验的验证方法。样本概览纹理版可分离卷积Description 与核心定位根据 README.md 的描述该样本是Texture-based implementation of a separable 2D convolution with a gaussian kernel. Used for performance comparison against convolutionSeparable.即基于纹理的高斯核二维可分离卷积实现用于与 convolutionSeparable 做性能对比。它的关键特性是完全不使用共享内存而是像 OpenGL 实现那样依赖纹理硬件路径完成邻域数据访问——这正是理解该样本的核心切入点。样本 main.cpp 的注释对此做了更直白的说明This sample implements the same algorithm as the convolutionSeparable CUDA Sample, but without using the shared memory at all. Instead, it uses textures in exactly the same way an OpenGL-based implementation would do.两个样本实现的是同一算法区别仅在数据通路一个是共享内存 寄存器缓存一个是纹理采样。这种同算法、不同实现的成对设计使开发者可以在同一张卡上直观对比两类访存模式的性能差异。Key Concepts关键概念Image Processing图像处理卷积是空间滤波的基础算子Texture纹理本样本的核心数据通路Data Parallel Algorithms数据并行算法每个输出像素由独立的线程计算。可分离卷积原理为什么能拆成两次一维卷积二维卷积直接计算时对每个输出像素需要遍历(2R1) × (2R1)个邻域输入。当卷积核可分离例如高斯核满足G(x,y) G(x) · G(y)时二维卷积可以分解为先做一次行方向一维卷积、再做一次列方向一维卷积计算量从O((2R1)²)降至O(2·(2R1))且两次卷积使用同一组一维核系数。本样本的 convolutionTexture_common.h 将核半径与核长度定义为编译期常量#define KERNEL_RADIUS 8 #define KERNEL_LENGTH (2 * KERNEL_RADIUS 1) // 即 17也就是说一维核长度为 17半径 8。对应的两条 GPU 路径分别是convolutionRowsKernel沿x 方向行做一维卷积输出到d_DstconvolutionColumnsKernel沿y 方向列做一维卷积。两次卷积之间需要把行卷积结果回写到纹理才能继续做列卷积——因为 CUDA 内核无法直接写入纹理详见下文主流程分析。从纹理引用到纹理对象API 演进与 CUDA API 清单本样本涉及的 CUDA Runtime APIREADME 明确列出本样本使用的 CUDA Runtime APIcudaMemcpy、cudaMallocArray、cudaFreeArray、cudaFree、cudaMemcpyToArray、cudaDeviceSynchronize、cudaCreateTextureObject、cudaMemcpyToSymbol、cudaMalloc其中cudaCreateTextureObject是最值得关注的一个它属于 CUDA 5.0 引入的纹理对象Texture Object机制将原先分散在texture引用声明、texRef.filterMode等绑定点上的信息统一收拢为cudaResourceDesc资源描述cudaTextureDesc采样描述 一个运行时创建的对象句柄。相较于纹理引用纹理对象可以直接作为内核函数参数传递这也是 convolutionRowsKernel 和 convolutionColumnsKernel 的签名中都带cudaTextureObject_t texSrc的原因。纹理对象的完整创建流程在 main.cpp 中可以看到标准三步式创建// 1. 分配 CUDA Array 作为纹理底层存储 cudaChannelFormatDesc floatTex cudaCreateChannelDescfloat(); checkCudaErrors(cudaMallocArray(a_Src, floatTex, imageW, imageH)); // 2. 填充资源描述声明纹理来自一个 CUDA Array cudaResourceDesc texRes; memset(texRes, 0, sizeof(cudaResourceDesc)); texRes.resType cudaResourceTypeArray; texRes.res.array.array a_Src; // 3. 填充采样描述并创建纹理对象 cudaTextureDesc texDescr; memset(texDescr, 0, sizeof(cudaTextureDesc)); texDescr.normalizedCoords false; // 使用像素坐标而非归一化坐标 texDescr.filterMode cudaFilterModeLinear; // 线性滤波 texDescr.addressMode[0] cudaAddressModeWrap; // x 方向环绕寻址 texDescr.addressMode[1] cudaAddressModeWrap; // y 方向环绕寻址 texDescr.readMode cudaReadModeElementType; checkCudaErrors(cudaCreateTextureObject(texSrc, texRes, texDescr, NULL));各采样参数的含义与影响参数本样本取值说明normalizedCoordsfalse纹理坐标使用整数像素坐标非归一化采样时以像素为单位偏移filterModecudaFilterModeLinear线性滤波模式配合tex2Dfloat返回浮点插值结果addressMode[0/1]cudaAddressModeWrap坐标越界时环绕wrap寻址与 OpenGL 的GL_REPEAT语义一致readModecudaReadModeElementType按元素类型直接读取不做归一化转换常量内存承载卷积核系数卷积核系数存放在常量内存中这是cudaMemcpyToSymbol的典型用法convolutionTexture.cu__constant__ float c_Kernel[KERNEL_LENGTH]; extern C void setConvolutionKernel(float *h_Kernel) { cudaMemcpyToSymbol(c_Kernel, h_Kernel, KERNEL_LENGTH * sizeof(float)); }常量内存的优势在于当一个 warp 内所有线程访问同一地址时硬件会将其广播为一次读取延迟极低。在卷积内层循环里所有线程读取的c_Kernel[i]完全相同因此这是教科书级的常量内存适用场景。核函数实现模板展开、像素坐标与线程网格模板递归展开Template-based Loop UnrollingconvolutionTexture.cu 使用 C 模板递归把内层 17 次卷积累加在编译期展开避免运行期循环开销template int i __device__ float convolutionRow(float x, float y, cudaTextureObject_t texSrc) { return tex2Dfloat(texSrc, x (float)(KERNEL_RADIUS - i), y) * c_Kernel[i] convolutionRowi - 1(x, y, texSrc); } template __device__ float convolutionRow-1(float x, float y, cudaTextureObject_t texSrc) { return 0; }convolutionColumn与之对称只是把采样偏移放在 y 方向。#define UNROLL_INNER 1convolutionTexture.cu控制走展开路径若置 0 则退化为运行期for (int k -KERNEL_RADIUS; k KERNEL_RADIUS; k)的普通循环convolutionTexture.cu方便对比展开与否的性能差异。线程网格与像素坐标映射两个核函数共享相同的网格配置convolutionTexture.cudim3 threads(16, 12); // 每个线程块 192 个线程 dim3 blocks(iDivUp(imageW, threads.x), iDivUp(imageH, threads.y));线程到像素的映射使用了__mul24优化的IMAD宏convolutionTexture.cu// Maps to a single instruction on G8x / G9x / G10x #define IMAD(a, b, c) (__mul24((a), (b)) (c))每个线程计算一个输出像素采样坐标取像素中心0.5f越界线程直接返回const float x (float)ix 0.5f; const float y (float)iy 0.5f; if (ix imageW || iy imageH) return;注意tex2Dfloat返回的是模板化浮点采样结果叠加线性滤波模式后卷积求和本质上是在做硬件加速的加权邻域聚合——这正是纹理路径相对共享内存路径的差异所在。内核无法直写纹理行→列之间的一次拷贝CUDA 内核不能直接写纹理因此行卷积结果必须先拷贝到普通设备内存d_Output再经cudaMemcpyToArray回填到纹理数组才能作为列卷积的输入。主流程中这一步被单独计时main.cpp其注释明确写道 While CUDA kernels cant write to textures directly, this copy is inevitable。checkCudaErrors( cudaMemcpyToArray(a_Src, 0, 0, d_Output, imageW * imageH * sizeof(float), cudaMemcpyDeviceToDevice));这是纹理方案相对共享内存方案的一个固有代价也是两个样本性能对比时值得关注的环节。主流程与性能测量3072×1536 图像上的基准测试数据规模与初始化main.cpp 固定了测试规模const int imageW 3072; const int imageH 3072 / 2; // 1536 const unsigned int iterations 10;输入图像与核系数由srand(2009)固定种子的随机数生成rand() % 16保证结果可复现main.cpp。设备选择走findCudaDevice(argc, argv)来自 Common/helper_cuda.h默认选算力最高的设备。计时策略使用sdkCreateTimer/sdkResetTimer/sdkStartTimer/sdkStopTimer来自 Common/helper_functions.h关键点是在计时前调用cudaDeviceSynchronize()确保前序异步操作完成后再计时行卷积循环iterations次取平均耗时并换算为吞吐量imageW * imageH * 1e-6 / (0.001 * gpuTime)单位为Mpix/s行→纹理拷贝单次计时列卷积同样循环 10 次取平均。输出形如Average convolutionRowsGPU() time: %f msecs; //%f Mpix/s cudaMemcpyToArray() time: %f msecs; //%f Mpix/s Average convolutionColumnsGPU() time: %f msecs; //%f Mpix/s这与 doc/Performance.xls 记录的对比数据相互印证该表即为本样本与 convolutionSeparable 在多种架构上的性能对照。正确性验证CPU 参考实现与相对 L2 范数CPU 参考卷积convolutionTexture_gold.cpp 提供两个串行参考函数convolutionRowsCPU/convolutionColumnsCPU与 GPU 版本的差异在于边界处理CPU 版对越界坐标做**钳位clamp**到边缘像素int d x k; if (d 0) d 0; if (d imageW) d imageW - 1;而 GPU 版依赖纹理的cudaAddressModeWrap环绕寻址。两种边界策略在图像内部区域结果一致边缘区域则由插值/寻址语义决定因此最终校验采用容差化的 L2 范数而非逐像素精确比对。相对 L2 范数判定主程序先跑 CPU 行、列卷积得到参考输出h_OutputCPU再与 GPU 输出h_OutputGPU计算相对 L2 范数main.cppdouble delta 0, sum 0; for (unsigned int i 0; i imageW * imageH; i) { sum h_OutputCPU[i] * h_OutputCPU[i]; delta (h_OutputGPU[i] - h_OutputCPU[i]) * (h_OutputGPU[i] - h_OutputCPU[i]); } double L2norm sqrt(delta / sum);判定阈值L2norm 1e-6则打印Test failed!并返回失败否则打印Test passedmain.cpp。这是一种工程上常用的浮点结果验证方式既容忍纹理硬件插值与钳位边界带来的微小数值差异又能严格约束整体误差上界。平台支持、构建与运行支持的平台与架构来自 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即 Maxwell 到 Hopper/Blackwell 一代的完整覆盖支持的操作系统Linux、Windows支持的 CPU 架构x86_64、armv7l。前提条件按 README 说明需先安装对应平台的 CUDA Toolkit 明确该版本支持 CUDA Toolkit 13.3。编译与运行样本的 CMakeLists.txt 声明了cxx_std_17 cuda_std_17、启用CUDA_SEPARABLE_COMPILATION默认编译架构为75 80 86 87 89 90 100 110 120并支持-DENABLE_CUDA_DEBUGTrue开启cuda-gdb调试对应-G编译选项。目标由三个源文件构成add_executable(convolutionTexture convolutionTexture.cu convolutionTexture_gold.cpp main.cpp)遵循仓库统一的构建流程可在仓库根目录或直接在该样本子目录执行mkdir build cd build cmake .. make -j$(nproc)随后运行生成的convolutionTexture可执行程序程序会自动选择设备、跑完行/列卷积与 CPU 校验并打印各阶段耗时与Test passed。与 convolutionSeparable 的对比纹理 vs 共享内存本样本的存在意义在于与兄弟样本做同算法性能对比二者文件构成一一对应维度convolutionTextureconvolutionSeparable数据通路纹理对象 tex2D采样convolutionTexture.cu__shared__共享内存 寄存器convolutionSeparable.cu边界处理纹理地址模式wrap显式钳位clamp额外代价行结果需cudaMemcpyToArray回写纹理需处理 halo边界冗余加载参考文档doc/Performance.xlsdoc/convolutionSeparable.pdf其 Performance 章节讨论了纹理方案的取舍从源码结构可以推断共享内存方案通过一次协作加载把 tile 数据放入s_Data再用 halo 处理边界而纹理方案把边界语义和邻域读取全部交给纹理硬件代码更接近 OpenGL 式的数据流。实际性能取决于具体 GPU 架构与纹理缓存行为这正是两份样本各自附带性能数据表、供开发者自行实测对比的原因。小结convolutionTexture 是一个小而完整的纹理编程教学样本串联了纹理对象创建cudaCreateTextureObject、CUDA Array 存储、常量内存cudaMemcpyToSymbol、tex2D采样、模板展开优化、CPU 参考验证与相对 L2 范数判定等一整套关键知识点。开发者既可以把它当作纹理 API 的入门模板也可以直接在本目录与 convolutionSeparable 之间切换编译、实测对比形成对共享内存 vs 纹理两条卷积数据通路性能差异的第一手认识。【免费下载链接】cuda-samplesSamples for CUDA Developers which demonstrates features in CUDA Toolkit项目地址: https://gitcode.com/GitHub_Trending/cu/cuda-samples创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考