代码之家  ›  专栏  ›  技术社区  ›  Athanasios Margaris

如何将CUDA线程与三重嵌套循环相关联?

  •  0
  • Athanasios Margaris  · 技术社区  · 2 年前

    假设有三个变量X、Y和Z分别具有最小、最大和变化步长值,即minX、maxX、minY、maxY、minZ、maxZ(最小值和最大值可以有任何值,负值或正值)和步长X、步长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,threadId x.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)的解决方案,我会将其适用于其他情况。

    1 回复  |  直到 2 年前
        1
  •  1
  •   Soonts    2 年前

    形式为(threadIdx.x,threadIdx.y,threadId x.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 );
    

    如果你想用运行C++的CPU来测试以上内容, #include <cmath> using std::fma;

    函数求值的总数xCircles x yCircles x zCircles,可能具有数十亿的值

    是的,生成数十亿个CUDA线程可能是次优的。

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

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