纹理内存加速二维图像插值的前提是访问具有二维空间局部性；错配访问模式反而更慢，因纹理缓存命中率低、硬件优化失效。
纹理内存对二维图像插值采样有明确加速效果，但前提是访问模式匹配其硬件优化逻辑——不是所有“用 tex2D”的操作都能变快，错配反而拖慢。
为什么 tex2D 在图像插值中比全局内存快
纹理单元的缓存专为二维空间局部性设计：当相邻线程（比如一个 16×16 线程块）读取图像上邻近坐标时，硬件会预取整块 2D 数据进纹理缓存（Texture Cache），而不是单个 float；而全局内存即使合并访问（coalesced），也只按 32/64 字节事务加载，对插值所需的跨像素邻域（如双线性需要 4 个纹素）无感知。
关键点在于：tex2D 的浮点坐标 + 硬件插值是原子操作，一次调用就完成地址计算、缓存查找、四点加权求和——这些在全局内存中得手写 4 次 [] 访问 + 手动插值，不仅代码长，还破坏 coalescing。
- 仅当线程访问坐标在二维平面上呈聚集分布（非随机跳转）时，纹理缓存命中率才高
-
cudaReadModeNormalizedFloat模式下，uchar4图像数据自动转[0.0,1.0]浮点，省去手动归一化除法 - 启用
cudaFilterModeLinear后，tex2D内部直接走双线性插值流水线，延迟远低于 kernel 中模拟
绑定二维纹理必须用 cudaBindTexture 还是 cudaCreateTextureObject
用 cudaCreateTextureObject。旧式 cudaBindTexture 接口已废弃（CUDA 11.0+ 编译警告），且不支持统一内存（Unified Memory）和 CUDA 图（Graph）；更重要的是，它要求显式管理绑定/解绑生命周期，容易在多 kernel 复用时出错。
正确做法是：先分配 CUDA 数组（cudaMalloc3DArray），把图像数据拷入其中，再用 cudaCreateTextureObject 创建对象，传入 kernel 作为参数：
cudaChannelFormatDesc desc = cudaCreateChannelDesc<float>();
cudaArray* cuArray;
cudaMalloc3DArray(&cuArray, &desc, make_cudaExtent(w, h, 0));
// 拷贝图像数据到 cuArray...
cudaResourceDesc resDesc = {};
resDesc.resType = cudaResourceTypeArray;
resDesc.res.array.array = cuArray;
cudaTextureDesc texDesc = {};
texDesc.readMode = cudaReadModeElementType;
texDesc.filterMode = cudaFilterModeLinear; // 开启双线性
texDesc.addressMode[0] = cudaAddressModeClamp;
texDesc.addressMode[1] = cudaAddressModeClamp;
cudaTextureObject_t texObj = 0;
cudaCreateTextureObject(&texObj, &resDesc, &texDesc, nullptr);
// 传入 kernel：kernel>>(texObj, ...);
</float>
- 必须用
cudaMalloc3DArray而非cudaMalloc分配线性内存再绑定——后者不启用纹理缓存的 2D 优化路径 -
addressMode设为cudaAddressModeClamp防止边界外采样越界崩溃，比在 kernel 里写min(max(x,0),w-1)更高效 - 若图像原始是
uchar4，改用cudaCreateChannelDesc<uchar4>()</uchar4>并设readMode = cudaReadModeNormalizedFloat，可一步到位输出float4
tex2D 坐标偏移 0.5f 是必须的吗
不是必须，但绝大多数图像处理场景下必须加。原因在于纹理坐标默认以纹素（texel）中心为原点：一个 w×h 的图像，合法整数坐标范围是 (0.5f, 0.5f) 到 (w-0.5f, h-0.5f)；如果直接用像素索引 (x, y) 调用 tex2D，会采样到两像素交界处，导致结果偏移半个像素。
典型修正方式：
- 缩放 kernel 中：输入坐标按比例映射后，统一加
+0.5f再传给tex2D - 旋转/仿射变换时：变换矩阵需预先补偿 0.5 偏移，否则插值中心错位
- 若你用的是归一化坐标（
[0.0,1.0]），则偏移应为+0.5f/w和+0.5f/h，而非固定0.5f
漏掉这个偏移，图像边缘会出现半像素撕裂或模糊不对称，调试时极难定位。
哪些情况用纹理内存反而更慢
纹理内存不是银弹。以下场景它会拖累性能：
- 线程访问完全随机：比如光线追踪中每条光线打到屏幕不同位置且无空间相关性，纹理缓存命中率趋近于 0，每次
tex2D都触发全局内存加载 - 单线程高频串行读：纹理单元吞吐针对并行访存优化，单个线程反复调用
tex2D不如直接读 global memory + 寄存器缓存 - 数据频繁更新：纹理对象绑定后，若图像内容在 host 端改了但没重新 upload 到
cudaArray，采样结果就是脏数据——而全局内存 memcpy 后立即可见 - 小尺寸图像（
）：纹理缓存块大小通常 ≥128 字节，小图无法填满缓存行，带宽利用率反不如 L1 缓存
最易被忽略的一点：纹理内存只加速读，不提供原子写或任何形式的写回。如果你的算法需要“读-改-写”循环，必须换回全局内存或共享内存。
C++免费学习笔记（深入）：立即使用
在学习笔记中，你将探索 C++ 的入门与实战技巧！
