（cudaBindTexture2D）如何从中间绑定一个倾斜的数组

Posted 2023-03-23

【中文标题】（cudaBindTexture2D）如何从中间绑定一个倾斜的数组

【英文标题】：(cudaBindTexture2D) How to bind a pitched-array from the middle

【发布时间】：2022-01-09 00:06:35

【问题描述】：

我正在尝试从中间部分绑定一个倾斜的数组（而不是从数组的开头），如下所示。

/1.分配/

 cudaMallocPitch((void**)&d_texinput, &FloatPitch, cols*sizeof(float), rows);
 cudaMallocPitch((void**)&d_output, &FloatPitch, cols*sizeof(float), rows);

/2.设置目标区域的行长（即分行10次）/

 int row_div_times = 10;
 int part_rows = rows / row_div_times;
 int part_offset = part_rows*FloatPitch/sizeof(float);
 dim3 threads(16,16);
 dim3 Part_Blocks((cols + threads.x - 1) / threads.x, (Part_rows + threads.y - 1) / threads.y);

/3.迭代处理分割的行/

 for (int i = 0; i < row_div_times; i++)
 
      size_t offsetsize= i*part_offset;            

      /*computing values of "d_tex_input"*/
      calibration << <Part_Blocks, threads, 0, stream[i] >> >
         (d_texinput + i*part_offset );         

      /*
      //###(QUESTION point!) I want to bind the device memory "d_texinput" to texture "tex_mem" only partly like below.
      cudaBindTexture2D(0, tex_mem, &d_texinput[i*part_offset], channelDesc_flt, cols, Part_rows, FloatPitch); //tentative code a;
      ,,, or something like,,,
      cudaBindTexture2D(&offsetsize, tex_mem, &d_texinput, channelDesc_flt, cols, Part_rows, FloatPitch); //tentative code b;
      */
       


       //final computaion with texture
         final_computationwithtexture << <Part_Blocks, threads, 0, stream[i] >> >
         ( d_output + i*part_offset ); 

       cudaUnbindTexture(tex_mem);

 

请允许我请教您，建议如何通过修改以上部分来绑定设备内存阵列的目标区域（问题点！）？

我试图将 cudaBindTExture2D 的第一个参数理解为“偏移量”。但这不是价值。它是地址。根据文档。 我仍然无法理解文档。 我希望我能通过了解 cudaBindTexture2D 的适当输入方式来理解那是什么。

【参考方案1】：

offset 参数不是输入，而是输出。这就是为什么它是一个指针。该函数将以字节为单位设置偏移量。如果您想在分配过程中进行绑定，请适当设置 devPtr 参数（第三个），然后该函数将为您提供纹理访问所需的偏移量。

这是如何理解的：纹理只能以一定的对齐方式绑定。内存分配总是正确对齐的。因此，在大多数情况下，这不是问题。但是，如果您提供任意内存地址，CUDA 必须向下舍入到对齐，并且您必须稍后应用适当的偏移量。

假设您绑定了&float[66]，正确的对齐方式可能是&float[64]，因此CUDA 从该偏移量开始其纹理，您必须为每次访问添加8 个字节的偏移量才能获得所需的结果。我这里是选随机数，不知道对齐要求。

【讨论】：

非常感谢您的及时指导和很好的例子。因为我在音高的每个开头都读取了数组，听起来这个移位问题不会是问题，但了解这种基本行为对我来说非常重要。我明白了，现在我知道第一个参数“偏移量”的含义以及为什么需要清楚地输出它。非常感谢！

@mush_head 是的，在代码中添加assert(offset == 0) 可能是个好主意，以防万一，但应该没问题。