CUDA 11.3中已废弃的tex2D<T=float>重载接口的替代及使用方法

触发该编译警告的接口，是CUDA早期基于编译期全局静态绑定的老式纹理采样重载，在CUDA 11.3版本中被正式标记为废弃。其官方指定替代方案是基于运行时纹理对象（cudaTextureObject_t）的纹理采样体系。

替代接口说明 #

被标记废弃的重载签名为tex2D(texture<T, 2, cudaReadModeElementType>, float, float)，对应的无废弃警告的同名重载接口参数结构如下：

// 设备端核函数内的调用形式
SampleType val = tex2D(cudaTextureObject_t texObj, float x, float y);

该接口的采样精度、滤波逻辑、边界处理规则和旧版接口完全兼容，仅替换了纹理句柄的传递形式，不会改变原有采样的计算结果。

落地使用步骤 #

• 移除原代码中全局域定义的编译期绑定texture变量，不再依赖静态纹理绑定逻辑
• 在主机端核函数启动前，完成两个关键描述符的配置：
  • 资源描述符cudaResourceDesc：指定待采样显存的指针、数据类型、宽度、高度、行步长等基础内存布局信息
  • 纹理描述符cudaTextureDesc：指定坐标归一化模式、边界采样规则、滤波类型、读取模式等采样逻辑参数
• 调用cudaCreateTextureObject()接口，基于上述两个描述符生成可跨核函数传递的cudaTextureObject_t句柄
• 核函数形参替换为接收cudaTextureObject_t类型的纹理对象，函数内部将原来传入旧texture变量的采样调用，替换为传入该纹理对象句柄即可

新旧实现对照 #

触发废弃警告的旧写法

// 全局域静态绑定的纹理变量
__device__ texture<float, 2, cudaReadModeElementType> static_tex;

__global__ void old_sample_kernel() {
    // 该行会触发目标废弃警告
    float sample_val = tex2D(static_tex, x_coord, y_coord);
}

无废弃警告的新写法

// 主机端核函数启动前的纹理对象创建逻辑
cudaResourceDesc res_desc = {};
res_desc.resType = cudaResourceTypePitch2D;
res_desc.res.pitch2D.devPtr = d_device_buffer;
res_desc.res.pitch2D.desc = cudaCreateChannelDesc<float>();
res_desc.res.pitch2D.width = image_width;
res_desc.res.pitch2D.height = image_height;
res_desc.res.pitch2D.pitchInBytes = buffer_pitch;

cudaTextureDesc tex_desc = {};
tex_desc.addressMode[0] = cudaAddressModeClamp;
tex_desc.addressMode[1] = cudaAddressModeClamp;
tex_desc.filterMode = cudaFilterModeLinear;
tex_desc.readMode = cudaReadModeElementType;
tex_desc.normalizedCoords = false;

cudaTextureObject_t new_tex_obj = 0;
cudaCreateTextureObject(&new_tex_obj, &res_desc, &tex_desc, nullptr);

// 核函数定义
__global__ void new_sample_kernel(cudaTextureObject_t tex_obj) {
    // 无废弃警告的采样调用
    float sample_val = tex2D(tex_obj, x_coord, y_coord);
}

// 启动核函数时传入纹理对象句柄
new_sample_kernel<<<grid_dim, block_dim>>>(new_tex_obj);

过渡兼容说明 #

如果短期没有精力完成全量代码改造，可以通过给nvcc添加编译参数-D_CUDA_NO_TEXTURE_OBJECT_DEPRECATION_WARNING临时屏蔽该类废弃警告，但该方案仅为短期过渡手段，未来CUDA正式版本会完全移除旧版静态纹理接口，长期开发仍需迁移到纹理对象体系。
核函数执行完成后，可通过cudaDestroyTextureObject()销毁不再使用的纹理对象，释放对应的硬件资源。

内容的提问来源于stack exchange，提问作者Soleil