企业官网建设流程全解析

CUDA Samples bindlessTexture 示例深度解析：cudaTextureObject / cudaSurfaceObject 与 MipMap 的无绑定纹理实践

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

本文基于 CUDA Samples 仓库（当前支持 CUDA Toolkit 13.3）中的 cpp/3_CUDA_Features/bindlessTexture 示例，讲解 CUDA 无绑定（bindless）纹理体系的核心用法：如何通过cudaTextureObject与cudaSurfaceObject消除传统纹理 API 的全局绑定限制，如何用一张"图集纹理（atlas）"存放 64 位纹理对象句柄实现虚拟纹理（virtual texturing），以及如何在设备端直接以内核参数传递纹理/表面对象并程序化生成 MipMap。读完本文，你将掌握无绑定纹理对象从创建、传递到销毁的完整生命周期，并能将该模式复用到稀疏纹理、多纹理采样等实际渲染场景中。

示例概览：它到底演示了什么

根据 bindlessTexture/README.md 的说明，该示例的核心目标是演示cudaSurfaceObject、cudaTextureObject以及 CUDA 中的 MipMap 支持。运行该示例需要Compute Capability SM 3.0 及以上的 GPU——这一限制的根源在于无绑定纹理对象（cudaTextureObject_t）是 CUDA 5.0 引入的特性，底层依赖 Kepler 及以后架构的设备端句柄支持。

示例的文件结构如下：

从 内核文件头部注释 可以确认整体设计思路：示例包含两个内核——一个负责每帧渲染，另一个负责在启动阶段生成 MipMap 各级别；渲染采用"虚拟纹理"方案，即用一张 2D 纹理存放对真实纹理的引用，这正是通过 CUDA 5.0 引入的cudaTextureObject实现的。

无绑定纹理：为什么需要它

在 CUDA 5.0 之前，纹理通过__texture__等全局绑定点（binding point）使用，纹理与内核之间的绑定关系是全局且静态的，内核内无法灵活切换纹理、无法用数组管理纹理集合。无绑定纹理对象（bindless texture object）则把"纹理资源的描述"封装成一个 64 位的cudaTextureObject_t句柄，该句柄可以：

• 像普通变量一样作为内核参数传递；
• 存储在全局内存、常量内存乃至其他纹理中；
• 由cudaCreateTextureObject在运行时动态创建、由cudaDestroyTextureObject销毁。

类似地，cudaSurfaceObject_t允许内核通过surf2Dwrite等 API 直接写入表面对象，无需全局绑定点，因此可以像传参数一样把"输出表面"传入内核。

本示例正好把这两个特性组合成了完整的渲染管线：图集纹理存句柄 → 内核解码句柄 → 采样真实纹理 → 写入表面对象生成 MipMap。

虚拟纹理：用图集存放纹理对象句柄

设计思想

示例用一张4×4的 2D 图集纹理（atlas）来模拟虚拟纹理页表。图集的每个纹素不再存颜色，而是存一个64 位的cudaTextureObject_t句柄。渲染内核先采样图集拿到句柄，再解码句柄去采样对应的真实纹理，从而在一帧里动态决定"这一块应该显示哪张纹理"。

核心的编解码函数位于 bindlessTexture_kernel.cu：

__host__ __device__ __inline__ uint2 encodeTextureObject(cudaTextureObject_t obj) { return make_uint2((uint)(obj & 0xFFFFFFFF), (uint)(obj >> 32)); } __host__ __device__ __inline__ cudaTextureObject_t decodeTextureObject(uint2 obj) { return (((cudaTextureObject_t)obj.x) | ((cudaTextureObject_t)obj.y) << 32); }

注意uint2本身是 64 位，恰好可以容纳一个cudaTextureObject_t句柄，因此图集纹理用cudaCreateChannelDesc<uint2>()描述通道格式（见 initAtlasAndImages）。

渲染内核

渲染内核 d_render 的完整流程是：

__global__ void d_render(uchar4 *d_output, uint imageW, uint imageH, float lod, cudaTextureObject_t atlasTexture) { uint x = blockIdx.x * blockDim.x + threadIdx.x; uint y = blockIdx.y * blockDim.y + threadIdx.y; float u = x / (float)imageW; float v = y / (float)imageH; if ((x < imageW) && (y < imageH)) { // 从 2D 图集纹理中读出编码后的纹理对象 uint2 texCoded = tex2D<uint2>(atlasTexture, u, v); cudaTextureObject_t tex = decodeTextureObject(texCoded); // 用 tex2DLod 直接指定 mip map 层级采样 float4 color = tex2DLod<float4>(tex, u, 1 - v, lod); uint i = y * imageW + x; d_output[i] = to_uchar4(color * 255.0); } }

值得注意的几个技术点：

1. 模板化的tex2D<uint2>：图集纹理被声明为uint2元素类型，内核通过模板参数指定返回数据类型，这与 CUDA 5.0 引入的通用纹理读取函数配套使用。
2. tex2DLod<float4>：该函数允许直接传入显式的 LOD（mip map 层级）值，这正是交互控制 MipMap 层级的入口。内核注释还提到同批 API 中的tex2DGrad可以通过传递导数实现自动 MipMap/各向异性过滤。
3. u/v坐标：图集纹理与真实纹理都采用归一化坐标（normalizedCoords = 1），其中 v 方向做了1 - v翻转以匹配纹理坐标约定。

随机化图集

randomizeAtlas 为图集每个纹素随机分配三张内容纹理（flower、person、sponge）之一的句柄，然后通过cudaMemcpy3D配合make_cudaPitchedPtr把主机端uint2数据拷贝进图集数组。渲染时按下r键即可触发重新随机化，直观看到虚拟纹理"页表重映射"的效果。

MipMap：内核级生成与采样

创建 MipMap 数组

在 initAtlasAndImages 中，每张内容纹理都按如下步骤建立完整 MipMap 链：

1. 用 getMipMapLevels 计算所需层级数（取长宽深最大值反复除以 2 直到 0）；
2. cudaMallocMipmappedArray分配uchar4格式的 mipmapped array；
3. cudaGetMipmappedArrayLevel拿到第 0 级cudaArray_t，用cudaMemcpy3D上传原始图像；
4. 调用generateMipMaps逐级生成更高层级。

最终生成的纹理对象描述（bindlessTexture_kernel.cu）集中体现了 MipMap 采样的关键配置：

cudaResourceDesc resDescr; memset(&resDescr, 0, sizeof(cudaResourceDesc)); resDescr.resType = cudaResourceTypeMipmappedArray; resDescr.res.mipmap.mipmap = image.mipmapArray; cudaTextureDesc texDescr; memset(&texDescr, 0, sizeof(cudaTextureDesc)); texDescr.normalizedCoords = 1; // 归一化坐标 texDescr.filterMode = cudaFilterModeLinear; // 空间过滤 texDescr.mipmapFilterMode = cudaFilterModeLinear; // mip 层间线性插值 texDescr.addressMode[0..2] = cudaAddressModeClamp; // 寻址模式 texDescr.maxMipmapLevelClamp = float(levels - 1); // 限制最大层级 texDescr.readMode = cudaReadModeNormalizedFloat; checkCudaErrors(cudaCreateTextureObject(&image.textureObject, &resDescr, &texDescr, NULL));

用 Surface Object 生成 MipMap

MipMap 生成内核 d_mipmap 演示了无绑定表面对象的典型用法——它同时接收一个cudaSurfaceObject_t（输出）和cudaTextureObject_t（输入）作为内核参数：

__global__ void d_mipmap(cudaSurfaceObject_t mipOutput, cudaTextureObject_t mipInput, uint imageW, uint imageH) { uint x = blockIdx.x * blockDim.x + threadIdx.x; uint y = blockIdx.y * blockDim.y + threadIdx.y; float px = 1.0 / float(imageW); float py = 1.0 / float(imageH); if ((x < imageW) && (y < imageH)) { // 取 4 个像素的平均，生成高一层的 mip float4 color = (tex2D<float4>(mipInput, (x + 0) * px, (y + 0) * py)) + (tex2D<float4>(mipInput, (x + 1) * px, (y + 0) * py)) + (tex2D<float4>(mipInput, (x + 1) * px, (y + 1) * py)) + (tex2D<float4>(mipInput, (x + 0) * px, (y + 1) * py)); color /= 4.0; color *= 255.0; color = fminf(color, make_float4(255.0)); surf2Dwrite(to_uchar4(color), mipOutput, x * sizeof(uchar4), y); } }

生成流程 generateMipMaps 在主机端循环执行：

• 每次循环用cudaGetMipmappedArrayLevel取相邻两级数组（levelFrom/levelTo）；
• 用cudaArrayGetInfo校验目标级尺寸符合预期（checkHost断言）；
• 为levelFrom创建临时纹理对象（cudaResourceTypeArray，线性过滤、Clamp 寻址、归一化坐标）；
• 为levelTo创建临时表面对象；
• 以16×16块启动d_mipmap，写完后cudaDestroySurfaceObject/cudaDestroyTextureObject及时释放临时对象。

注释特别指出：使用归一化坐标访问（(x+0)*px这种形式）是为了保证非 2 的幂次纹理在下采样时行为正确。这正是 Surface Object 相对传统绑定表面（bound surface）的关键优势：表面不再占用全局绑定槽，可以按需创建、用完即毁。

CUDA 与 OpenGL 互操作：PBO 渲染管线

示例同时演示了 "Graphics Interop" 这一关键概念（README 的 Key Concepts 一节将其与 Texture 并列）。主机端通过 GLUT 创建 512×512 窗口，渲染结果经 CUDA 内核写入 OpenGL 像素缓冲对象（PBO），再由 OpenGL 绘制上屏，形成"CUDA 计算 → OpenGL 显示"的流水线。

缓冲区注册与映射

在 initGLBuffers 中，PBO 被注册为 CUDA 图形资源：

glGenBuffers(1, &pbo); glBindBuffer(GL_PIXEL_UNPACK_BUFFER_ARB, pbo); glBufferData(GL_PIXEL_UNPACK_BUFFER_ARB, windowSize.x * windowSize.y * sizeof(GLubyte) * 4, 0, GL_STREAM_DRAW_ARB); glBindBuffer(GL_PIXEL_UNPACK_BUFFER_ARB, 0); checkCudaErrors(cudaGraphicsGLRegisterBuffer(&cuda_pbo_resource, pbo, cudaGraphicsMapFlagsWriteDiscard));

每帧渲染时（render 函数）依次调用：

• cudaGraphicsMapResources映射资源；
• cudaGraphicsResourceGetMappedPointer取得设备指针d_output；
• 启动renderAtlasImage内核；
• cudaGraphicsUnmapResources解除映射，之后 OpenGL 才能安全读取 PBO。

display回调随后用glBindBuffer(GL_PIXEL_UNPACK_BUFFER_ARB, pbo)+glDrawPixels把 PBO 内容绘制到窗口，并使用sdkGetAverageTimerValue统计 FPS 写入窗口标题。程序退出时，cleanup调用cudaGraphicsUnregisterResource注销图形资源并删除 PBO。

交互操作

程序启动后打印如下操作提示（bindlessTexture.cpp）：

空闲回调 idle 在动画开启时每帧让lod += 0.02，配合renderAtlasImage中对 LOD 的三角波折叠（fmodf+fabs，见 bindlessTexture_kernel.cu），实现 MipMap 层级在 0 到最高层之间往复扫描的视觉效果。

构建与运行

环境依赖

README 的 Dependencies 一节列出的构建/运行依赖为 X11、OpenGL、Freeglut、GLEW，这些依赖的详细说明见仓库根 README.md。此外必须安装与平台匹配的 CUDA Toolkit。从 CMakeLists.txt 可以看到，若系统未找到 OpenGL 或 GLUT，该示例会在构建时自动跳过（输出 "GLUT not found" / "OpenGL not found" 提示），这符合仓库"依赖缺失即 waive 自身"的构建策略。

CMake 构建

示例的 CMakeLists.txt 声明LANGUAGES C CXX CUDA，默认目标架构列表为75 80 86 87 89 90 100 110 120，编译选项包括--extended-lambda、cxx_std_17/cuda_std_17及CUDA_SEPARABLE_COMPILATION ON。既可以在仓库根目录统一构建，也可以进入示例子目录单独构建：

# 在仓库根目录 mkdir build && cd build cmake .. make -j$(nproc) # 或单独构建本示例 cd cpp/3_CUDA_Features/bindlessTexture mkdir build && cd build cmake .. make -j$(nproc)

Windows 下需使用 Visual Studio 提供的x64 Native Tools Command Prompt for VS，可用cmake -G "Visual Studio 16 2019" -A x64配置。构建完成后在 build 目录对应子目录下运行./bindlessTexture（Linux）即可，data/目录中的 PPM 贴图会通过 POST_BUILD 命令自动拷贝到输出目录。

自动化验证模式

该示例支持无窗口的自动验证：传入-file参数指向参考数据文件后，程序跳过 OpenGL 初始化，直接渲染一帧并与参考输出逐像素比对。仓库的 test_args.json 中配置了标准测试参数：

"bindlessTexture": { "args": [ "-file=data/ref_bindlessTexture.bin" ] }

验证流程在 runAutoTest 中实现：cudaMalloc输出缓冲 → 启动renderAtlasImage→cudaDeviceSynchronize→cudaMemcpy回主机 →sdkDumpBin导出bindlessTexture.bin→ 用sdkCompareBin2BinFloat与参考文件比对（误差阈值MAX_EPSILON_ERROR = 5.0f、THRESHOLD = 0.15f）。这也意味着该示例可以纳入仓库根目录的python3 run_tests.py --output ./test --dir ./build/cpp --config test_args.json批量回归测试。

生命周期管理：从创建到销毁

无绑定纹理对象属于显式资源，必须成对管理。示例给出了完整的资源清理范式：

• 创建：cudaMallocArray/cudaMallocMipmappedArray分配存储，cudaCreateTextureObject/cudaCreateSurfaceObject创建句柄；
• 销毁：deinitAtlasAndImages 依次释放每张内容纹理的主机数据、cudaDestroyTextureObject、cudaFreeMipmappedArray，再释放图集的主机数据、纹理对象与cudaFreeArray；
• 主机端：cleanup注销 CUDA 图形资源并删除 GL 缓冲，通过 GLUT 的glutCloseFunc或atexit注册到退出路径。

此外，临时创建的 MipMap 输入纹理对象与输出表面对象在 generateMipMaps 每次迭代后立即销毁，避免句柄泄漏。开发者在自己的无绑定纹理代码中应遵循同样的对称生命周期原则，并配合checkCudaErrors/getLastCudaError对每次 API 调用与内核启动做错误检查。

小结

bindlessTexture 示例用约 400 行代码把 CUDA 无绑定纹理体系的三个支柱串联成一个可交互、可验证的完整程序：

1. cudaTextureObject作参数传递：句柄可存于纹理（图集虚拟纹理）、可传内核、可在设备端解码；
2. cudaSurfaceObject作写入目标：MipMap 生成内核通过surf2Dwrite直接写表面，摆脱全局绑定点的束缚；
3. MipMap 全链路支持：从cudaMallocMipmappedArray、cudaGetMipmappedArrayLevel到tex2DLod显式 LOD 采样与maxMipmapLevelClamp配置。

对需要实现稀疏纹理、大规模纹理集合管理或程序化 mip 生成的开发者而言，bindlessTexture_kernel.cu 中的模式（atlas 存句柄 + 内核级编解码 + 临时对象按需创建销毁）是可直接复用的参考实现；而 bindlessTexture.cpp 则提供了 CUDA 计算与 OpenGL 显示互操作的完整样板。

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