如何将CUDA线程与三重嵌套循环关联起来？

Question

假设有三个变量 X、Y 和 Z，分别具有最小、最大和变化步长值，minX、maxX、minY、maxY、minZ、maxZ（最小和最大值可以为任何值，负数或正数）和 stepX 、步骤Y和步骤Z，分别。我们还定义一个函数 W = f(X,Y,Z)，它将这些变量作为参数并返回值 W。现在，我们要估计定义的 X、Y 和 Z 的所有可能组合的 W 值通过上述参数。可以说，基于上述信息，我们已经估计了相关嵌套结构中三个循环中每一个的迭代次数 xCircles、yCircles 和 zCircles（来自 minX、maxX、minY、maxY、minZ、maxZ）。在这种情况下，该过程的串行版本的实现如下。

double curX, curY, curZ, W;
curX = minX; curY = minY; curZ = minZ;
for (long i=0;i<xCircles;i++) {
     for (long j=0;j<yCircles;j++) {
          curZ=minZ;
          for (long k=0;k<zCircles;k++) {
               W = f(curX,curY,curZ);
               curZ = minZ+(k+1)*stepZ;
          curY = minY+(j+1)*stepY; }
     curX = minX+(i+1)*stepX; }}        

目标是并行化这个串行三重循环。通过 #pragma omp parallel 和相关指令，使用 OpenMP 进行并行化非常容易，现在我想在 CUDA 中做同样的事情。我认为关于每个块的线程数和网格中的块数以及相关维度的最佳内核执行配置是一个实验问题（尽管一些技巧将受到赞赏），并且我致力于构建threadIdx 与三重循环每次迭代中使用的坐标 curX、curY 和 curZ 之间的映射，即 (threadIdx.x, threadIdx.y, threadIdx.z) ---> (curX, curY, curZ) 形式的映射。这样，我们就可以为每个线程分配f(threadIdx.x, threadIdx.y, threadIdx.z)的估计值，以完全消除三重循环。请记住，对于 minX、maxX、minY、maxY、minZ、maxZ 的大值以及 stepX、stepY 和 stepZ 的小值，函数计算 xCircles x yCircles x zCircles 的总数可能具有数十亿的值，并且它是可以合理地假设每个线程不仅执行一次，而且执行大量计算，其数量可以被视为问题的参数。有关于如何做到这一点的任何想法吗？或者，有人知道有此类问题的书或论文可以开始阅读吗？事实上，这个问题有四种不同类型的嵌套循环，深度为 d=1,2,3,4，但如果我找到针对这种特殊情况的解决方案 (d=3)，我会将其适应其他情况。

Answer 1

(threadIdx.x, threadIdx.y, threadIdx.z) 形式的映射 ---> (curX, curY, curZ)

从整数索引计算 curX、curY、curZ 似乎很简单，具体方法如下：

double curX = fma( i, stepX, minX );
double curY = fma( j, stepY, minY );
double curZ = fma( k, stepZ, minZ );

如果您想使用 CPU 运行的 C++ 来测试上述内容，

#include <cmath>

和

using std::fma;

函数评估总数xCircles x yCircles x zCircles，可能有数十亿的价值

是的，生成数十亿个 CUDA 线程可能不是最理想的。

一个典型的解决方法是为某个外部循环的每次迭代生成 CUDA 线程，但在 CUDA 内核内部编写内部循环，即在那里编写一个

for

循环。

例如，当 xCircles=yCircles=zCircles=1000 时，您可以启动 1000x1000 个 CUDA 线程的网格，并在内核中编写运行 1000 次迭代的内部循环。