如何使用cudaBindTexture2D绑定pitched数组的中间分段区域

错误原因分析 #

• 写法a错误：
  1. 第三个参数错误加了取地址符&，&d_texinput[i*part_offset]得到的是主机端存储指针变量的地址，不是设备端分段的起始地址
  2. 第一个参数传0不符合API要求，该参数是输出型参数，需要传入size_t类型变量的地址接收返回的字节偏移
• 写法b错误：
  第三个参数传了整个数组的起始地址d_texinput，没有指向目标分段的起始位置，即使得到偏移也需要额外在核函数中处理坐标，不符合你的需求

正确绑定代码 #

你需要的绑定分段的正确写法如下，直接替换注释中的错误写法即可：

size_t tex_offset;
// 绑定当前分段到纹理
cudaBindTexture2D(&tex_offset, tex_mem, d_texinput + i * part_offset, channelDesc_flt, cols, part_rows, FloatPitch);

关键说明 #

1. 第一个参数&tex_offset是API的输出参数，会返回绑定地址相对于纹理对齐要求的字节偏移，通过cudaMallocPitch分配的地址天然满足对齐要求，多数场景下该值为0，无需额外处理
2. 第三个参数直接传偏移后的设备指针即可，不要加取地址符：d_texinput本身是设备端float类型指针，i * part_offset是按float单位计算的偏移量，相加后刚好是目标分段的设备起始地址
3. 绑定分段时行步长FloatPitch保持原值即可，因为pitched数组的全局行对齐属性不会随分段改变
4. 核函数中访问纹理时，坐标(0, 0)直接对应当前绑定分段的第一个元素，不需要额外累加全局行偏移

更优实现方案（推荐） #

你可以省略循环内的重复绑定/解绑操作，直接在循环外绑定整个数组，在核函数访问纹理时手动叠加行偏移即可，能减少API调用开销：

// 循环外一次性绑定整个数组
size_t tex_offset;
cudaBindTexture2D(&tex_offset, tex_mem, d_texinput, channelDesc_flt, cols, rows, FloatPitch);

for (int i = 0; i < row_div_times; i++)
{
    size_t offsetsize= i*part_offset;
    calibration <<<Part_Blocks, threads, 0, stream[i]>>>(d_texinput + i*part_offset );

    // 核函数传入当前分段的行偏移量，访问纹理时y坐标叠加该值即可
    final_computationwithtexture <<<Part_Blocks, threads, 0, stream[i]>>>(d_output + i*part_offset, i * part_rows);
}
// 循环外统一解绑
cudaUnbindTexture(tex_mem);

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