c - CUDA 中的图像序列处理

Question

我的任务是使用CUDA实现图像重建算法代码。我获得了相同的 C 代码。代码的输入是一个 DAT 文件，其中包含大小为 640 x 480 的 360 个图像。代码如下所示:

    FILE *in,*out;
    float *i_data,*o_data;
    i_data=(float *)malloc(mem_size);
    for(int projection=0;projection<360;projection++)
    {
      in=fopen("filename.dat","rb");
      fread(i_data,mem_size,1,in);
      ... some math ...
      for(int slice_no=-240;slice_no<240:slice_no++)
      {
          for (i=-320;i<320;i++)
          for (j=-320;j<320;j++)
          {
             // do some operations
             (*(o_data*slice_no)+(j+320)+(i+240))+=(*(i_data*value)+(j+240)+(i+320));
             // some more math
          }
      } 
    }

输出浮点指针被写回 dat 文件。如果我必须并行化这些循环，我将如何在 CUDA 中做到这一点？我尝试使用 640 个线程中的每个线程使用 640 个 block 来实现 CUDA 中的内部两个 for 循环。如何将线程索引赋予循环内的指针操作。我尝试给予

         int i=blockIdx.x;
         int j=threadIdx.x;

     and 

         kernel<<<640,640>>>

但这在输出指针中给出了错误的值。大多数是NAN。除了上面代码片段中显示的带有指针的行之外，我能够成功地实现其他数学。

有人可以帮我做这件事吗？我在这段代码中做错了什么？另外，是否可以并行化这里的所有 for 循环？

Answer 1

将内部循环限制在单个 block 内运行效率不高。此外，如果并行化所有三个循环(或者如果您可以从磁盘加载更多数据并一次将许多图像发送到 GPU，则并行化 4 个循环)，您将获得更多好处。

你的问题看起来像

for(x = minx; x < maxx; ++x)
{
  for(y = miny; y < maxy; ++y)
  {
    for(z = minz; z < maxz; ++z)
    {
      // do some math
    }
   }
}

内核中的线程数应该是:

num_thread = (maxx-minx) * (maxy-miny) * (maxz-minz);

你应该独立于 block 的大小。将其设置为像 blockDim.x = 256 这样的常量(通过实验找到更好的常量)。

区 block 数量为

gridDim.x =(num_thread + blockDim.x) / blockDim.x

请注意，我的推理就好像这三个循环展开为单个大循环，但如果您的驱动程序允许，您可以制作 2D 甚至 3D 线程 block 。

在您的内核中，您可以根据内置变量计算大展开循环内的全局索引，如下所示

int index = blockIdx.x * blockDim.x + threadIdx.x

确保您没有超出范围

if(index >= num_thread) return; // do nothing

现在使用index，您可以恢复x、y和z(slice_no, i, j) 为

x = index / ( (maxy-miny) * (maxz-minz) ) + minx;
y = ( index % ( (maxy-miny) * (maxz-minz) ) ) / (maxz-minz) + miny;
z = ( index % ( (maxy-miny) * (maxz-minz) ) ) % (maxz-minz) + minz;