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 及以后架构的设备端句柄支持。示例的文件结构如下文件职责bindlessTexture.cpp主机端入口GLUT 窗口、OpenGL PBO 互操作、键盘交互、自动验证bindlessTexture.h定义Image数据结构图像尺寸、数组句柄、纹理对象句柄bindlessTexture_kernel.cu设备端核心代码虚拟纹理渲染内核、MipMap 生成内核、纹理/图集初始化data/三张 PPM 贴图flower、person、sponge与自动化验证参考数据 ref_bindlessTexture.binCMakeLists.txtCMake 构建脚本依赖 OpenGL 与 GLUT从 内核文件头部注释 可以确认整体设计思路示例包含两个内核——一个负责每帧渲染另一个负责在启动阶段生成 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句柄因此图集纹理用cudaCreateChannelDescuint2()描述通道格式见 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 tex2Duint2(atlasTexture, u, v); cudaTextureObject_t tex decodeTextureObject(texCoded); // 用 tex2DLod 直接指定 mip map 层级采样 float4 color tex2DLodfloat4(tex, u, 1 - v, lod); uint i y * imageW x; d_output[i] to_uchar4(color * 255.0); } }值得注意的几个技术点模板化的tex2Duint2图集纹理被声明为uint2元素类型内核通过模板参数指定返回数据类型这与 CUDA 5.0 引入的通用纹理读取函数配套使用。tex2DLodfloat4该函数允许直接传入显式的 LODmip map 层级值这正是交互控制 MipMap 层级的入口。内核注释还提到同批 API 中的tex2DGrad可以通过传递导数实现自动 MipMap/各向异性过滤。u/v坐标图集纹理与真实纹理都采用归一化坐标normalizedCoords 1其中 v 方向做了1 - v翻转以匹配纹理坐标约定。随机化图集randomizeAtlas 为图集每个纹素随机分配三张内容纹理flower、person、sponge之一的句柄然后通过cudaMemcpy3D配合make_cudaPitchedPtr把主机端uint2数据拷贝进图集数组。渲染时按下r键即可触发重新随机化直观看到虚拟纹理页表重映射的效果。MipMap内核级生成与采样创建 MipMap 数组在 initAtlasAndImages 中每张内容纹理都按如下步骤建立完整 MipMap 链用 getMipMapLevels 计算所需层级数取长宽深最大值反复除以 2 直到 0cudaMallocMipmappedArray分配uchar4格式的 mipmapped arraycudaGetMipmappedArrayLevel拿到第 0 级cudaArray_t用cudaMemcpy3D上传原始图像调用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 生成 MipMapMipMap 生成内核 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 (tex2Dfloat4(mipInput, (x 0) * px, (y 0) * py)) (tex2Dfloat4(mipInput, (x 1) * px, (y 0) * py)) (tex2Dfloat4(mipInput, (x 1) * px, (y 1) * py)) (tex2Dfloat4(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及时释放临时对象。注释特别指出使用归一化坐标访问(x0)*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按键功能空格切换自动动画animate并将 LOD 重置为 0/LOD 增加 0.25-LOD 减少 0.25r随机化虚拟图集重新分配各纹素指向的纹理句柄Esc退出程序空闲回调 idle 在动画开启时每帧让lod 0.02配合renderAtlasImage中对 LOD 的三角波折叠fmodffabs见 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 目录对应子目录下运行./bindlessTextureLinux即可data/目录中的 PPM 贴图会通过 POST_BUILD 命令自动拷贝到输出目录。自动化验证模式该示例支持无窗口的自动验证传入-file参数指向参考数据文件后程序跳过 OpenGL 初始化直接渲染一帧并与参考输出逐像素比对。仓库的 test_args.json 中配置了标准测试参数bindlessTexture: { args: [ -filedata/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 无绑定纹理体系的三个支柱串联成一个可交互、可验证的完整程序cudaTextureObject作参数传递句柄可存于纹理图集虚拟纹理、可传内核、可在设备端解码cudaSurfaceObject作写入目标MipMap 生成内核通过surf2Dwrite直接写表面摆脱全局绑定点的束缚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创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考
CUDA Samples bindlessTexture 示例深度解析：cudaTextureObject / cudaSurfaceObject 与 MipMap 的无绑定纹理实践
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 及以后架构的设备端句柄支持。示例的文件结构如下文件职责bindlessTexture.cpp主机端入口GLUT 窗口、OpenGL PBO 互操作、键盘交互、自动验证bindlessTexture.h定义Image数据结构图像尺寸、数组句柄、纹理对象句柄bindlessTexture_kernel.cu设备端核心代码虚拟纹理渲染内核、MipMap 生成内核、纹理/图集初始化data/三张 PPM 贴图flower、person、sponge与自动化验证参考数据 ref_bindlessTexture.binCMakeLists.txtCMake 构建脚本依赖 OpenGL 与 GLUT从 内核文件头部注释 可以确认整体设计思路示例包含两个内核——一个负责每帧渲染另一个负责在启动阶段生成 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句柄因此图集纹理用cudaCreateChannelDescuint2()描述通道格式见 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 tex2Duint2(atlasTexture, u, v); cudaTextureObject_t tex decodeTextureObject(texCoded); // 用 tex2DLod 直接指定 mip map 层级采样 float4 color tex2DLodfloat4(tex, u, 1 - v, lod); uint i y * imageW x; d_output[i] to_uchar4(color * 255.0); } }值得注意的几个技术点模板化的tex2Duint2图集纹理被声明为uint2元素类型内核通过模板参数指定返回数据类型这与 CUDA 5.0 引入的通用纹理读取函数配套使用。tex2DLodfloat4该函数允许直接传入显式的 LODmip map 层级值这正是交互控制 MipMap 层级的入口。内核注释还提到同批 API 中的tex2DGrad可以通过传递导数实现自动 MipMap/各向异性过滤。u/v坐标图集纹理与真实纹理都采用归一化坐标normalizedCoords 1其中 v 方向做了1 - v翻转以匹配纹理坐标约定。随机化图集randomizeAtlas 为图集每个纹素随机分配三张内容纹理flower、person、sponge之一的句柄然后通过cudaMemcpy3D配合make_cudaPitchedPtr把主机端uint2数据拷贝进图集数组。渲染时按下r键即可触发重新随机化直观看到虚拟纹理页表重映射的效果。MipMap内核级生成与采样创建 MipMap 数组在 initAtlasAndImages 中每张内容纹理都按如下步骤建立完整 MipMap 链用 getMipMapLevels 计算所需层级数取长宽深最大值反复除以 2 直到 0cudaMallocMipmappedArray分配uchar4格式的 mipmapped arraycudaGetMipmappedArrayLevel拿到第 0 级cudaArray_t用cudaMemcpy3D上传原始图像调用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 生成 MipMapMipMap 生成内核 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 (tex2Dfloat4(mipInput, (x 0) * px, (y 0) * py)) (tex2Dfloat4(mipInput, (x 1) * px, (y 0) * py)) (tex2Dfloat4(mipInput, (x 1) * px, (y 1) * py)) (tex2Dfloat4(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及时释放临时对象。注释特别指出使用归一化坐标访问(x0)*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按键功能空格切换自动动画animate并将 LOD 重置为 0/LOD 增加 0.25-LOD 减少 0.25r随机化虚拟图集重新分配各纹素指向的纹理句柄Esc退出程序空闲回调 idle 在动画开启时每帧让lod 0.02配合renderAtlasImage中对 LOD 的三角波折叠fmodffabs见 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 目录对应子目录下运行./bindlessTextureLinux即可data/目录中的 PPM 贴图会通过 POST_BUILD 命令自动拷贝到输出目录。自动化验证模式该示例支持无窗口的自动验证传入-file参数指向参考数据文件后程序跳过 OpenGL 初始化直接渲染一帧并与参考输出逐像素比对。仓库的 test_args.json 中配置了标准测试参数bindlessTexture: { args: [ -filedata/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 无绑定纹理体系的三个支柱串联成一个可交互、可验证的完整程序cudaTextureObject作参数传递句柄可存于纹理图集虚拟纹理、可传内核、可在设备端解码cudaSurfaceObject作写入目标MipMap 生成内核通过surf2Dwrite直接写表面摆脱全局绑定点的束缚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创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考